while-conversion.mlir

// *** IR Dump After Canonicalizer ***


module {
  func @while() attributes {iree.module.export} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    %2 = "mhlo.while"(%0) ( {
    ^bb0(%arg0: tensor<i32>):  // no predecessors
      %3 = "mhlo.compare"(%arg0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
      "mhlo.return"(%3) : (tensor<i1>) -> ()
    },  {
    ^bb0(%arg0: tensor<i32>):  // no predecessors
      %3 = mhlo.add %arg0, %arg0 : tensor<i32>
      "mhlo.return"(%3) : (tensor<i32>) -> ()
    }) : (tensor<i32>) -> tensor<i32>
    check.expect_eq(%2, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::mhlo::(anonymous namespace)::LegalizeControlFlowPass ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::(anonymous namespace)::HLOToHLOPreprocessing ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After RemoveShapeConstraints ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::ConvertShapeToShapex ***


module {
  func @while() attributes {iree.module.export} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %4 = extract_element %3[] : tensor<i1>
    cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%5: tensor<i32>):  // pred: ^bb1
    %6 = mhlo.add %5, %5 : tensor<i32>
    br ^bb1(%6 : tensor<i32>)
  ^bb3(%7: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%7, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::(anonymous namespace)::FlattenTuplesInCFGPass ***


module {
  func @while() attributes {iree.module.export} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %4 = extract_element %3[] : tensor<i1>
    cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%5: tensor<i32>):  // pred: ^bb1
    %6 = mhlo.add %5, %5 : tensor<i32>
    br ^bb1(%6 : tensor<i32>)
  ^bb3(%7: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%7, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After Inliner ***


module {
  func @while() attributes {iree.module.export} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %4 = extract_element %3[] : tensor<i1>
    cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%5: tensor<i32>):  // pred: ^bb1
    %6 = mhlo.add %5, %5 : tensor<i32>
    br ^bb1(%6 : tensor<i32>)
  ^bb3(%7: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%7, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::LegalizeInputTypesPass ***


module {
  func @while() attributes {iree.module.export} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %4 = extract_element %3[] : tensor<i1>
    cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%5: tensor<i32>):  // pred: ^bb1
    %6 = mhlo.add %5, %5 : tensor<i32>
    br ^bb1(%6 : tensor<i32>)
  ^bb3(%7: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%7, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::MaterializeExportedReflectionPass ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::ExpandFunctionDynamicDimsPass ***
func @while() attributes {iree.module.export} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::MergeExportedReflectionPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::TieDynamicShapesPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::MaterializeShapeCalculationsPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::HoistShapeCalculations ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = extract_element %3[] : tensor<i1>
  cond_br %4, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%5: tensor<i32>):  // pred: ^bb1
  %6 = mhlo.add %5, %5 : tensor<i32>
  br ^bb1(%6 : tensor<i32>)
^bb3(%7: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%7, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::PrePartitioningConversionPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %4 = "mhlo.convert"(%3) : (tensor<i1>) -> tensor<i8>
  %5 = flow.tensor.load %4 : tensor<i8>
  %6 = trunci %5 : i8 to i1
  cond_br %6, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%7: tensor<i32>):  // pred: ^bb1
  %8 = mhlo.add %7, %7 : tensor<i32>
  br ^bb1(%8 : tensor<i32>)
^bb3(%9: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%9, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::DispatchabilityAnalysisPass ***


module {
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = "mhlo.compare"(%2, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %4 = "mhlo.convert"(%3) : (tensor<i1>) -> tensor<i8>
    %5 = flow.tensor.load %4 : tensor<i8>
    %6 = trunci %5 : i8 to i1
    cond_br %6, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%7: tensor<i32>):  // pred: ^bb1
    %8 = mhlo.add %7, %7 : tensor<i32>
    br ^bb1(%8 : tensor<i32>)
  ^bb3(%9: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%9, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::(anonymous namespace)::IdentifyDispatchRegions2Pass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %c1 = constant 1 : index
  %c1_2 = constant 1 : index
  %3 = flow.dispatch.region[%c1_2 : index](%arg0 = %2 : tensor<i32>, %arg1 = %1 : tensor<i32>) -> tensor<i1> {
    %10 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    flow.return %10 : tensor<i1>
  }
  %4 = flow.dispatch.region[%c1 : index](%arg0 = %3 : tensor<i1>) -> tensor<i8> {
    %10 = "mhlo.convert"(%arg0) : (tensor<i1>) -> tensor<i8>
    flow.return %10 : tensor<i8>
  }
  %5 = flow.tensor.load %4 : tensor<i8>
  %6 = trunci %5 : i8 to i1
  cond_br %6, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%7: tensor<i32>):  // pred: ^bb1
  %c1_3 = constant 1 : index
  %8 = flow.dispatch.region[%c1_3 : index](%arg0 = %7 : tensor<i32>) -> tensor<i32> {
    %10 = mhlo.add %arg0, %arg0 : tensor<i32>
    flow.return %10 : tensor<i32>
  }
  br ^bb1(%8 : tensor<i32>)
^bb3(%9: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%9, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %c1 = constant 1 : index
  %3 = flow.dispatch.region[%c1 : index](%arg0 = %2 : tensor<i32>, %arg1 = %1 : tensor<i32>) -> tensor<i1> {
    %10 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    flow.return %10 : tensor<i1>
  }
  %4 = flow.dispatch.region[%c1 : index](%arg0 = %3 : tensor<i1>) -> tensor<i8> {
    %10 = "mhlo.convert"(%arg0) : (tensor<i1>) -> tensor<i8>
    flow.return %10 : tensor<i8>
  }
  %5 = flow.tensor.load %4 : tensor<i8>
  %6 = trunci %5 : i8 to i1
  cond_br %6, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%7: tensor<i32>):  // pred: ^bb1
  %8 = flow.dispatch.region[%c1 : index](%arg0 = %7 : tensor<i32>) -> tensor<i32> {
    %10 = mhlo.add %arg0, %arg0 : tensor<i32>
    flow.return %10 : tensor<i32>
  }
  br ^bb1(%8 : tensor<i32>)
^bb3(%9: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%9, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::FoldCompatibleDispatchRegionsPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %c1 = constant 1 : index
  %3 = flow.dispatch.region[%c1 : index](%arg0 = %2 : tensor<i32>, %arg1 = %1 : tensor<i32>) -> tensor<i8> {
    %9 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %10 = "mhlo.convert"(%9) : (tensor<i1>) -> tensor<i8>
    flow.return %10 : tensor<i8>
  }
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch.region[%c1 : index](%arg0 = %6 : tensor<i32>) -> tensor<i32> {
    %9 = mhlo.add %arg0, %arg0 : tensor<i32>
    flow.return %9 : tensor<i32>
  }
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::RematerializeDispatchConstantsPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %c1 = constant 1 : index
  %3 = flow.dispatch.region[%c1 : index](%arg0 = %2 : tensor<i32>, %arg1 = %1 : tensor<i32>) -> tensor<i8> {
    %9 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %10 = "mhlo.convert"(%9) : (tensor<i1>) -> tensor<i8>
    flow.return %10 : tensor<i8>
  }
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch.region[%c1 : index](%arg0 = %6 : tensor<i32>) -> tensor<i32> {
    %9 = mhlo.add %arg0, %arg0 : tensor<i32>
    flow.return %9 : tensor<i32>
  }
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::OutlineDispatchRegionsPass ***


module {
  flow.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_0
    module {
      func @while_ex_dispatch_0(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i8> {
        %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
        %1 = "mhlo.convert"(%0) : (tensor<i1>) -> tensor<i8>
        return %1 : tensor<i8>
      }
    }
  }
  flow.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_1
    module {
      func @while_ex_dispatch_1(%arg0: tensor<i32>) -> tensor<i32> {
        %0 = mhlo.add %arg0, %arg0 : tensor<i32>
        return %0 : tensor<i32>
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %cst = constant dense<1> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<4> : tensor<i32>
    %0 = iree.do_not_optimize(%cst) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %c1 = constant 1 : index
    %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
    %4 = flow.tensor.load %3 : tensor<i8>
    %5 = trunci %4 : i8 to i1
    cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%6: tensor<i32>):  // pred: ^bb1
    %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
    br ^bb1(%7 : tensor<i32>)
  ^bb3(%8: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%8, %cst_1) : tensor<i32>
    return
  }
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %c1 = constant 1 : index
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::PostPartitioningConversionPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %c1 = constant 1 : index
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %c1 = constant 1 : index
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %cst = constant dense<1> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<4> : tensor<i32>
  %c1 = constant 1 : index
  %0 = iree.do_not_optimize(%cst) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst_1) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::(anonymous namespace)::HoistUnstreamableOps ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %cst = constant dense<4> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<1> : tensor<i32>
  %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst) : tensor<i32>
  return
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %cst = constant dense<4> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<1> : tensor<i32>
  %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%c1 : index](%2, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%c1 : index](%6) : (tensor<i32>) -> tensor<i32>
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst) : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::Flow::FormStreamsPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %cst = constant dense<4> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<1> : tensor<i32>
  %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
    %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
    flow.return %9 : tensor<i8>
  }
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
    %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
    flow.return %9 : tensor<i32>
  }
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst) : tensor<i32>
  return
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %cst = constant dense<4> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<1> : tensor<i32>
  %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
    %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
    flow.return %9 : tensor<i8>
  }
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
    %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
    flow.return %9 : tensor<i32>
  }
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst) : tensor<i32>
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %cst = constant dense<4> : tensor<i32>
  %cst_0 = constant dense<3> : tensor<i32>
  %cst_1 = constant dense<1> : tensor<i32>
  %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
  %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
  br ^bb1(%0 : tensor<i32>)
^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
  %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
    %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
    flow.return %9 : tensor<i8>
  }
  %4 = flow.tensor.load %3 : tensor<i8>
  %5 = trunci %4 : i8 to i1
  cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
^bb2(%6: tensor<i32>):  // pred: ^bb1
  %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
    %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
    flow.return %9 : tensor<i32>
  }
  br ^bb1(%7 : tensor<i32>)
^bb3(%8: tensor<i32>):  // pred: ^bb1
  check.expect_eq(%8, %cst) : tensor<i32>
  return
}

// *** IR Dump After SymbolDCE ***


module {
  flow.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_0
    module {
      func @while_ex_dispatch_0(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i8> {
        %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
        %1 = "mhlo.convert"(%0) : (tensor<i1>) -> tensor<i8>
        return %1 : tensor<i8>
      }
    }
  }
  flow.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_1
    module {
      func @while_ex_dispatch_1(%arg0: tensor<i32>) -> tensor<i32> {
        %0 = mhlo.add %arg0, %arg0 : tensor<i32>
        return %0 : tensor<i32>
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %c1 = constant 1 : index
    %cst = constant dense<4> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<1> : tensor<i32>
    %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
      %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
      flow.return %9 : tensor<i8>
    }
    %4 = flow.tensor.load %3 : tensor<i8>
    %5 = trunci %4 : i8 to i1
    cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%6: tensor<i32>):  // pred: ^bb1
    %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
      %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
      flow.return %9 : tensor<i32>
    }
    br ^bb1(%7 : tensor<i32>)
  ^bb3(%8: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%8, %cst) : tensor<i32>
    return
  }
}

// *** IR Dump After Canonicalizer ***


module {
  flow.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_0
    module {
      func @while_ex_dispatch_0(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i8> {
        %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
        %1 = "mhlo.convert"(%0) : (tensor<i1>) -> tensor<i8>
        return %1 : tensor<i8>
      }
    }
  }
  flow.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    flow.dispatch.entry @while_ex_dispatch_1
    module {
      func @while_ex_dispatch_1(%arg0: tensor<i32>) -> tensor<i32> {
        %0 = mhlo.add %arg0, %arg0 : tensor<i32>
        return %0 : tensor<i32>
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %c1 = constant 1 : index
    %cst = constant dense<4> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<1> : tensor<i32>
    %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
      %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
      flow.return %9 : tensor<i8>
    }
    %4 = flow.tensor.load %3 : tensor<i8>
    %5 = trunci %4 : i8 to i1
    cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%6: tensor<i32>):  // pred: ^bb1
    %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
      %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
      flow.return %9 : tensor<i32>
    }
    br ^bb1(%7 : tensor<i32>)
  ^bb3(%8: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%8, %cst) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::MaterializeInterfacesPass ***


module {
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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 @while_ex_dispatch_0() {
          %c0 = constant 0 : index
          %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
          %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
          %2 = call @while_ex_dispatch_0_impl(%0, %1) : (tensor<i32>, tensor<i32>) -> tensor<i8>
          hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<i8>
          return
        }
        func @while_ex_dispatch_0_impl(%arg0: tensor<i32>, %arg1: tensor<i32>) -> tensor<i8> attributes {sym_visibility = "private"} {
          %0 = "mhlo.compare"(%arg0, %arg1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
          %1 = "mhlo.convert"(%0) : (tensor<i1>) -> tensor<i8>
          return %1 : tensor<i8>
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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 @while_ex_dispatch_1() {
          %c0 = constant 0 : index
          %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
          %1 = call @while_ex_dispatch_1_impl(%0) : (tensor<i32>) -> tensor<i32>
          hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
          return
        }
        func @while_ex_dispatch_1_impl(%arg0: tensor<i32>) -> tensor<i32> attributes {sym_visibility = "private"} {
          %0 = mhlo.add %arg0, %arg0 : tensor<i32>
          return %0 : tensor<i32>
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %c1 = constant 1 : index
    %cst = constant dense<4> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<1> : tensor<i32>
    %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
      %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
      flow.return %9 : tensor<i8>
    }
    %4 = flow.tensor.load %3 : tensor<i8>
    %5 = trunci %4 : i8 to i1
    cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%6: tensor<i32>):  // pred: ^bb1
    %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
      %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
      flow.return %9 : tensor<i32>
    }
    br ^bb1(%7 : tensor<i32>)
  ^bb3(%8: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%8, %cst) : tensor<i32>
    return
  }
}

// *** IR Dump After Inliner ***
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 @while_ex_dispatch_0() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
    %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
    %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
    hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::TieDynamicShapesPass ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::MaterializeShapeCalculationsPass ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::HoistShapeCalculations ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::DecomposeHLOClampPass ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToCompatibleHLOPass ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = "mhlo.compare"(%0, %1) {comparison_direction = "LT"} : (tensor<i32>, tensor<i32>) -> tensor<i1>
  %3 = "mhlo.convert"(%2) : (tensor<i1>) -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToLinalgOnTensorsPass ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %1 {
  ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
    %4 = cmpi "slt", %arg0, %arg1 : i32
    linalg.yield %4 : i1
  }: tensor<i32>, tensor<i32> -> tensor<i1>
  %3 = linalg.generic {args_in = 1 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %2 {
  ^bb0(%arg0: i1):  // no predecessors
    %4 = zexti %arg0 : i1 to i8
    linalg.yield %4 : i8
  }: tensor<i1> -> tensor<i8>
  hal.interface.store.tensor %3, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After LinalgFusionOfTensorOps ***
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 @while_ex_dispatch_0() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
    %2 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %1 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: tensor<i32>, tensor<i32> -> tensor<i8>
    hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After LinalgFoldUnitExtentDims ***
func @while_ex_dispatch_0() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
  %2 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %1 {
  ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
    %3 = cmpi "slt", %arg0, %arg1 : i32
    %4 = zexti %3 : i1 to i8
    linalg.yield %4 : i8
  }: tensor<i32>, tensor<i32> -> tensor<i8>
  hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<i8>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_0() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
    %2 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %1 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: tensor<i32>, tensor<i32> -> tensor<i8>
    hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After LinalgFusionOfTensorOps ***
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 @while_ex_dispatch_0() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = hal.interface.load.tensor @legacy_io::@arg1, offset = %c0 : tensor<i32>
    %2 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %1 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: tensor<i32>, tensor<i32> -> tensor<i8>
    hal.interface.store.tensor %2, @legacy_io::@ret0, offset = %c0 : tensor<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToLinalgOnBuffersPass ***
func @while_ex_dispatch_0() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
  %c0 = constant 0 : index
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
    %3 = cmpi "slt", %arg0, %arg1 : i32
    %4 = zexti %3 : i1 to i8
    linalg.yield %4 : i8
  }: memref<i32>, memref<i32>, memref<i8>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_0() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: memref<i32>, memref<i32>, memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_0() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: memref<i32>, memref<i32>, memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::LinalgTileAndFusePass ***
func @while_ex_dispatch_0() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
    %3 = cmpi "slt", %arg0, %arg1 : i32
    %4 = zexti %3 : i1 to i8
    linalg.yield %4 : i8
  }: memref<i32>, memref<i32>, memref<i8>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::SplitDispatchFunctionPass ***
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 @while_ex_dispatch_0() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: memref<i32>, memref<i32>, memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::LinalgTileAndFusePass ***
func @while_ex_dispatch_0() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
    %3 = cmpi "slt", %arg0, %arg1 : i32
    %4 = zexti %3 : i1 to i8
    linalg.yield %4 : i8
  }: memref<i32>, memref<i32>, memref<i8>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_0() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %2, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i8):  // no predecessors
      %3 = cmpi "slt", %arg0, %arg1 : i32
      %4 = zexti %3 : i1 to i8
      linalg.yield %4 : i8
    }: memref<i32>, memref<i32>, memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertToGPUPass ***
func @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
  %3 = load %1[] : memref<i32>
  %4 = load %2[] : memref<i32>
  %5 = load %0[] : memref<i8>
  %6 = cmpi "slt", %3, %4 : i32
  %7 = zexti %6 : i1 to i8
  store %7, %0[] : memref<i8>
  return
}

// *** IR Dump After ConvertAffineToStandard ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = load %0[] : memref<i8>
    %6 = cmpi "slt", %3, %4 : i32
    %7 = zexti %6 : i1 to i8
    store %7, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = cmpi "slt", %3, %4 : i32
    %6 = zexti %5 : i1 to i8
    store %6, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = cmpi "slt", %3, %4 : i32
    %6 = zexti %5 : i1 to i8
    store %6, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ResolveShapeOpsPass ***
func @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
  %3 = load %1[] : memref<i32>
  %4 = load %2[] : memref<i32>
  %5 = cmpi "slt", %3, %4 : i32
  %6 = zexti %5 : i1 to i8
  store %6, %0[] : memref<i8>
  return
}

// *** IR Dump After LegalizeStandardForSPIRV ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = cmpi "slt", %3, %4 : i32
    %6 = zexti %5 : i1 to i8
    store %6, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = cmpi "slt", %3, %4 : i32
    %6 = zexti %5 : i1 to i8
    store %6, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_0() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i8>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg1} : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %2[] : memref<i32>
    %5 = cmpi "slt", %3, %4 : i32
    %6 = zexti %5 : i1 to i8
    store %6, %0[] : memref<i8>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertToSPIRVPass ***
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>}>} {
  spv.module Logical GLSL450 {
    spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    spv.func @while_ex_dispatch_0() "None" attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
      %0 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
      %1 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
      %2 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
      %3 = spv.constant 0 : i32
      %4 = spv.AccessChain %1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      %5 = spv.Load "StorageBuffer" %4 : i32
      %6 = spv.constant 0 : i32
      %7 = spv.AccessChain %2[%6, %6] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      %8 = spv.Load "StorageBuffer" %7 : i32
      %9 = spv.SLessThan %5, %8 : i32
      %10 = spv.constant 0 : i32
      %11 = spv.constant 1 : i32
      %12 = spv.Select %9, %11, %10 : i1, i32
      %13 = spv.constant 0 : i32
      %14 = spv.constant 4 : i32
      %15 = spv.constant 8 : i32
      %16 = spv.SMod %13, %14 : i32
      %17 = spv.IMul %16, %15 : i32
      %18 = spv.constant 255 : i32
      %19 = spv.ShiftLeftLogical %18, %17 : i32, i32
      %20 = spv.Not %19 : i32
      %21 = spv.BitwiseAnd %12, %18 : i32
      %22 = spv.ShiftLeftLogical %21, %17 : i32, i32
      %23 = spv.constant 4 : i32
      %24 = spv.SDiv %13, %23 : i32
      %25 = spv.AccessChain %0[%13, %24] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      %26 = spv.AtomicAnd "Device" "AcquireRelease" %25, %20 : !spv.ptr<i32, StorageBuffer>
      %27 = spv.AtomicOr "Device" "AcquireRelease" %25, %22 : !spv.ptr<i32, StorageBuffer>
      spv.Return
    }
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After SPIRVLowerABIAttributes ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %1 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %2 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %3 = spv.constant 0 : i32
    %4 = spv.AccessChain %1[%3, %3] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %5 = spv.Load "StorageBuffer" %4 : i32
    %6 = spv.constant 0 : i32
    %7 = spv.AccessChain %2[%6, %6] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %8 = spv.Load "StorageBuffer" %7 : i32
    %9 = spv.SLessThan %5, %8 : i32
    %10 = spv.constant 0 : i32
    %11 = spv.constant 1 : i32
    %12 = spv.Select %9, %11, %10 : i1, i32
    %13 = spv.constant 0 : i32
    %14 = spv.constant 4 : i32
    %15 = spv.constant 8 : i32
    %16 = spv.SMod %13, %14 : i32
    %17 = spv.IMul %16, %15 : i32
    %18 = spv.constant 255 : i32
    %19 = spv.ShiftLeftLogical %18, %17 : i32, i32
    %20 = spv.Not %19 : i32
    %21 = spv.BitwiseAnd %12, %18 : i32
    %22 = spv.ShiftLeftLogical %21, %17 : i32, i32
    %23 = spv.constant 4 : i32
    %24 = spv.SDiv %13, %23 : i32
    %25 = spv.AccessChain %0[%13, %24] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %26 = spv.AtomicAnd "Device" "AcquireRelease" %25, %20 : !spv.ptr<i32, StorageBuffer>
    %27 = spv.AtomicOr "Device" "AcquireRelease" %25, %22 : !spv.ptr<i32, StorageBuffer>
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_0
  spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
}

// *** IR Dump After Canonicalizer ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 1 : i32
    %1 = spv.constant 0 : i32
    %2 = spv.constant 8 : i32
    %3 = spv.constant 255 : i32
    %4 = spv.constant 4 : i32
    %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %9 = spv.Load "StorageBuffer" %8 : i32
    %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %11 = spv.Load "StorageBuffer" %10 : i32
    %12 = spv.SLessThan %9, %11 : i32
    %13 = spv.Select %12, %0, %1 : i1, i32
    %14 = spv.SMod %1, %4 : i32
    %15 = spv.IMul %14, %2 : i32
    %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
    %17 = spv.Not %16 : i32
    %18 = spv.BitwiseAnd %13, %3 : i32
    %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
    %20 = spv.SDiv %1, %4 : i32
    %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
    %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_0
  spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
}

// *** IR Dump After CSE ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 1 : i32
    %1 = spv.constant 0 : i32
    %2 = spv.constant 8 : i32
    %3 = spv.constant 255 : i32
    %4 = spv.constant 4 : i32
    %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %9 = spv.Load "StorageBuffer" %8 : i32
    %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %11 = spv.Load "StorageBuffer" %10 : i32
    %12 = spv.SLessThan %9, %11 : i32
    %13 = spv.Select %12, %0, %1 : i1, i32
    %14 = spv.SMod %1, %4 : i32
    %15 = spv.IMul %14, %2 : i32
    %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
    %17 = spv.Not %16 : i32
    %18 = spv.BitwiseAnd %13, %3 : i32
    %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
    %20 = spv.SDiv %1, %4 : i32
    %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
    %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_0
  spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
}

// *** IR Dump After SPIRVUpdateVCE ***
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 1 : i32
    %1 = spv.constant 0 : i32
    %2 = spv.constant 8 : i32
    %3 = spv.constant 255 : i32
    %4 = spv.constant 4 : i32
    %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %9 = spv.Load "StorageBuffer" %8 : i32
    %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %11 = spv.Load "StorageBuffer" %10 : i32
    %12 = spv.SLessThan %9, %11 : i32
    %13 = spv.Select %12, %0, %1 : i1, i32
    %14 = spv.SMod %1, %4 : i32
    %15 = spv.IMul %14, %2 : i32
    %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
    %17 = spv.Not %16 : i32
    %18 = spv.BitwiseAnd %13, %3 : i32
    %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
    %20 = spv.SDiv %1, %4 : i32
    %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
    %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_0
  spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::TranslateExecutablesPass ***
hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
  hal.interface @legacy_io {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
  hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
  hal.executable.target "vulkan*" {
    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>}>} {
      spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
        spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
        spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
        spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
        spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
          %0 = spv.constant 1 : i32
          %1 = spv.constant 0 : i32
          %2 = spv.constant 8 : i32
          %3 = spv.constant 255 : i32
          %4 = spv.constant 4 : i32
          %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
          %9 = spv.Load "StorageBuffer" %8 : i32
          %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
          %11 = spv.Load "StorageBuffer" %10 : i32
          %12 = spv.SLessThan %9, %11 : i32
          %13 = spv.Select %12, %0, %1 : i1, i32
          %14 = spv.SMod %1, %4 : i32
          %15 = spv.IMul %14, %2 : i32
          %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
          %17 = spv.Not %16 : i32
          %18 = spv.BitwiseAnd %13, %3 : i32
          %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
          %20 = spv.SDiv %1, %4 : i32
          %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
          %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
          %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
          spv.Return
        }
        spv.EntryPoint "GLCompute" @while_ex_dispatch_0
        spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
      }
      hal.interface @legacy_io attributes {sym_visibility = "private"} {
        hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
        hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
        hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
      }
    }
  }
}

// *** IR Dump After Inliner ***
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 @while_ex_dispatch_1() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = mhlo.add %0, %0 : tensor<i32>
    hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::TieDynamicShapesPass ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = mhlo.add %0, %0 : tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::MaterializeShapeCalculationsPass ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = mhlo.add %0, %0 : tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::HoistShapeCalculations ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = mhlo.add %0, %0 : tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::DecomposeHLOClampPass ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = mhlo.add %0, %0 : tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToCompatibleHLOPass ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = mhlo.add %0, %0 : tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToLinalgOnTensorsPass ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %0 {
  ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
    %2 = addi %arg0, %arg1 : i32
    linalg.yield %2 : i32
  }: tensor<i32>, tensor<i32> -> tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After LinalgFusionOfTensorOps ***
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 @while_ex_dispatch_1() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %0 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: tensor<i32>, tensor<i32> -> tensor<i32>
    hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After LinalgFoldUnitExtentDims ***
func @while_ex_dispatch_1() {
  %c0 = constant 0 : index
  %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
  %1 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %0 {
  ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
    %2 = addi %arg0, %arg1 : i32
    linalg.yield %2 : i32
  }: tensor<i32>, tensor<i32> -> tensor<i32>
  hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_1() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %0 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: tensor<i32>, tensor<i32> -> tensor<i32>
    hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After LinalgFusionOfTensorOps ***
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 @while_ex_dispatch_1() {
    %c0 = constant 0 : index
    %0 = hal.interface.load.tensor @legacy_io::@arg0, offset = %c0 : tensor<i32>
    %1 = linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %0, %0 {
    ^bb0(%arg0: i32, %arg1: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: tensor<i32>, tensor<i32> -> tensor<i32>
    hal.interface.store.tensor %1, @legacy_io::@ret0, offset = %c0 : tensor<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertHLOToLinalgOnBuffersPass ***
func @while_ex_dispatch_1() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
  %c0 = constant 0 : index
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
    %2 = addi %arg0, %arg1 : i32
    linalg.yield %2 : i32
  }: memref<i32>, memref<i32>, memref<i32>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_1() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: memref<i32>, memref<i32>, memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_1() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: memref<i32>, memref<i32>, memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::LinalgTileAndFusePass ***
func @while_ex_dispatch_1() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
    %2 = addi %arg0, %arg1 : i32
    linalg.yield %2 : i32
  }: memref<i32>, memref<i32>, memref<i32>
  return
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::SplitDispatchFunctionPass ***
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 @while_ex_dispatch_1() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: memref<i32>, memref<i32>, memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::LinalgTileAndFusePass ***
func @while_ex_dispatch_1() {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
  ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
    %2 = addi %arg0, %arg1 : i32
    linalg.yield %2 : i32
  }: memref<i32>, memref<i32>, memref<i32>
  return
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_1() {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    linalg.generic {args_in = 2 : i64, args_out = 1 : i64, indexing_maps = [affine_map<() -> ()>, affine_map<() -> ()>, affine_map<() -> ()>], iterator_types = []} %1, %1, %0 {
    ^bb0(%arg0: i32, %arg1: i32, %arg2: i32):  // no predecessors
      %2 = addi %arg0, %arg1 : i32
      linalg.yield %2 : i32
    }: memref<i32>, memref<i32>, memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertToGPUPass ***
func @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = load %1[] : memref<i32>
  %3 = load %1[] : memref<i32>
  %4 = load %0[] : memref<i32>
  %5 = addi %2, %3 : i32
  store %5, %0[] : memref<i32>
  return
}

// *** IR Dump After ConvertAffineToStandard ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = load %0[] : memref<i32>
    %5 = addi %2, %3 : i32
    store %5, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = addi %2, %3 : i32
    store %4, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = addi %2, %3 : i32
    store %4, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ResolveShapeOpsPass ***
func @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
  %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
  %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
  %2 = load %1[] : memref<i32>
  %3 = load %1[] : memref<i32>
  %4 = addi %2, %3 : i32
  store %4, %0[] : memref<i32>
  return
}

// *** IR Dump After LegalizeStandardForSPIRV ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = addi %2, %3 : i32
    store %4, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After Canonicalizer ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = addi %2, %3 : i32
    store %4, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After CSE ***
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 @while_ex_dispatch_1() attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = iree.placeholder for "interface buffer" {binding = @legacy_io::@ret0} : memref<i32>
    %1 = iree.placeholder for "interface buffer" {binding = @legacy_io::@arg0} : memref<i32>
    %2 = load %1[] : memref<i32>
    %3 = load %1[] : memref<i32>
    %4 = addi %2, %3 : i32
    store %4, %0[] : memref<i32>
    return
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertToSPIRVPass ***
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>}>} {
  spv.module Logical GLSL450 {
    spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    spv.func @while_ex_dispatch_1() "None" attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}, vkspv.workgroup_count_from_result_shape = 1 : i32} {
      %0 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
      %1 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
      %2 = spv.constant 0 : i32
      %3 = spv.AccessChain %1[%2, %2] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      %4 = spv.Load "StorageBuffer" %3 : i32
      %5 = spv.constant 0 : i32
      %6 = spv.AccessChain %1[%5, %5] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      %7 = spv.Load "StorageBuffer" %6 : i32
      %8 = spv.IAdd %4, %7 : i32
      %9 = spv.constant 0 : i32
      %10 = spv.AccessChain %0[%9, %9] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
      spv.Store "StorageBuffer" %10, %8 : i32
      spv.Return
    }
  }
  hal.interface @legacy_io attributes {sym_visibility = "private"} {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
}

// *** IR Dump After SPIRVLowerABIAttributes ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %1 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %2 = spv.constant 0 : i32
    %3 = spv.AccessChain %1[%2, %2] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %4 = spv.Load "StorageBuffer" %3 : i32
    %5 = spv.constant 0 : i32
    %6 = spv.AccessChain %1[%5, %5] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %7 = spv.Load "StorageBuffer" %6 : i32
    %8 = spv.IAdd %4, %7 : i32
    %9 = spv.constant 0 : i32
    %10 = spv.AccessChain %0[%9, %9] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    spv.Store "StorageBuffer" %10, %8 : i32
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_1
  spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
}

// *** IR Dump After Canonicalizer ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 0 : i32
    %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %4 = spv.Load "StorageBuffer" %3 : i32
    %5 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %6 = spv.Load "StorageBuffer" %5 : i32
    %7 = spv.IAdd %4, %6 : i32
    %8 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    spv.Store "StorageBuffer" %8, %7 : i32
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_1
  spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
}

// *** IR Dump After CSE ***
spv.module Logical GLSL450 {
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 0 : i32
    %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %4 = spv.Load "StorageBuffer" %3 : i32
    %5 = spv.Load "StorageBuffer" %3 : i32
    %6 = spv.IAdd %4, %5 : i32
    %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    spv.Store "StorageBuffer" %7, %6 : i32
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_1
  spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
}

// *** IR Dump After SPIRVUpdateVCE ***
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
  spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
  spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
    %0 = spv.constant 0 : i32
    %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
    %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    %4 = spv.Load "StorageBuffer" %3 : i32
    %5 = spv.Load "StorageBuffer" %3 : i32
    %6 = spv.IAdd %4, %5 : i32
    %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
    spv.Store "StorageBuffer" %7, %6 : i32
    spv.Return
  }
  spv.EntryPoint "GLCompute" @while_ex_dispatch_1
  spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::TranslateExecutablesPass ***
hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
  hal.interface @legacy_io {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
  hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
  hal.executable.target "vulkan*" {
    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>}>} {
      spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
        spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
        spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
        spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
          %0 = spv.constant 0 : i32
          %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
          %4 = spv.Load "StorageBuffer" %3 : i32
          %5 = spv.Load "StorageBuffer" %3 : i32
          %6 = spv.IAdd %4, %5 : i32
          %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
          spv.Store "StorageBuffer" %7, %6 : i32
          spv.Return
        }
        spv.EntryPoint "GLCompute" @while_ex_dispatch_1
        spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
      }
      hal.interface @legacy_io attributes {sym_visibility = "private"} {
        hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
        hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
      }
    }
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::LinkExecutablesPass ***


module {
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 1 : i32
            %1 = spv.constant 0 : i32
            %2 = spv.constant 8 : i32
            %3 = spv.constant 255 : i32
            %4 = spv.constant 4 : i32
            %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %9 = spv.Load "StorageBuffer" %8 : i32
            %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %11 = spv.Load "StorageBuffer" %10 : i32
            %12 = spv.SLessThan %9, %11 : i32
            %13 = spv.Select %12, %0, %1 : i1, i32
            %14 = spv.SMod %1, %4 : i32
            %15 = spv.IMul %14, %2 : i32
            %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
            %17 = spv.Not %16 : i32
            %18 = spv.BitwiseAnd %13, %3 : i32
            %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
            %20 = spv.SDiv %1, %4 : i32
            %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
            %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_0
          spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 0 : i32
            %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %4 = spv.Load "StorageBuffer" %3 : i32
            %5 = spv.Load "StorageBuffer" %3 : i32
            %6 = spv.IAdd %4, %5 : i32
            %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            spv.Store "StorageBuffer" %7, %6 : i32
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_1
          spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %c1 = constant 1 : index
    %cst = constant dense<4> : tensor<i32>
    %cst_0 = constant dense<3> : tensor<i32>
    %cst_1 = constant dense<1> : tensor<i32>
    %0 = iree.do_not_optimize(%cst_1) : tensor<i32>
    %1 = iree.do_not_optimize(%cst_0) : tensor<i32>
    br ^bb1(%0 : tensor<i32>)
  ^bb1(%2: tensor<i32>):  // 2 preds: ^bb0, ^bb2
    %3 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %2 : tensor<i32>, %arg2 = %1 : tensor<i32>) -> tensor<i8> {
      %9 = flow.dispatch @while_ex_dispatch_0::@while_ex_dispatch_0[%arg0 : index](%arg1, %arg2) : (tensor<i32>, tensor<i32>) -> tensor<i8>
      flow.return %9 : tensor<i8>
    }
    %4 = flow.tensor.load %3 : tensor<i8>
    %5 = trunci %4 : i8 to i1
    cond_br %5, ^bb2(%2 : tensor<i32>), ^bb3(%2 : tensor<i32>)
  ^bb2(%6: tensor<i32>):  // pred: ^bb1
    %7 = flow.ex.stream.fragment(%arg0 = %c1 : index, %arg1 = %6 : tensor<i32>) -> tensor<i32> {
      %9 = flow.dispatch @while_ex_dispatch_1::@while_ex_dispatch_1[%arg0 : index](%arg1) : (tensor<i32>) -> tensor<i32>
      flow.return %9 : tensor<i32>
    }
    br ^bb1(%7 : tensor<i32>)
  ^bb3(%8: tensor<i32>):  // pred: ^bb1
    check.expect_eq(%8, %cst) : tensor<i32>
    return
  }
}

// *** IR Dump After mlir::iree_compiler::(anonymous namespace)::ConvertFlowToHALPass ***


module {
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 1 : i32
            %1 = spv.constant 0 : i32
            %2 = spv.constant 8 : i32
            %3 = spv.constant 255 : i32
            %4 = spv.constant 4 : i32
            %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %9 = spv.Load "StorageBuffer" %8 : i32
            %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %11 = spv.Load "StorageBuffer" %10 : i32
            %12 = spv.SLessThan %9, %11 : i32
            %13 = spv.Select %12, %0, %1 : i1, i32
            %14 = spv.SMod %1, %4 : i32
            %15 = spv.IMul %14, %2 : i32
            %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
            %17 = spv.Not %16 : i32
            %18 = spv.BitwiseAnd %13, %3 : i32
            %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
            %20 = spv.SDiv %1, %4 : i32
            %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
            %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_0
          spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 0 : i32
            %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %4 = spv.Load "StorageBuffer" %3 : i32
            %5 = spv.Load "StorageBuffer" %3 : i32
            %6 = spv.IAdd %4, %5 : i32
            %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            spv.Store "StorageBuffer" %7, %6 : i32
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_1
          spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
    %c1 = constant 1 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %dev_0 = hal.ex.shared_device : !hal.device
    %allocator_1 = hal.device.allocator %dev_0 : !hal.allocator
    %cbuffer_2 = hal.allocator.allocate.const %allocator_1, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_2 : !hal.buffer
    %dev_3 = hal.ex.shared_device : !hal.device
    %allocator_4 = hal.device.allocator %dev_3 : !hal.allocator
    %cbuffer_5 = hal.allocator.allocate.const %allocator_4, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_5 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_5) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_2) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
    %dev_6 = hal.ex.shared_device : !hal.device
    %allocator_7 = hal.device.allocator %dev_6 : !hal.allocator
    %sz = hal.allocator.compute_size %allocator_7, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator_7, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev_6, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %exe = hal.executable.lookup %dev_6, @while_ex_dispatch_0 : !hal.executable
    %executable_layout = hal.executable_layout.lookup %dev_6, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
    %c0 = constant 0 : index
    %allocator_8 = hal.buffer.allocator %2 : !hal.allocator
    %sz_9 = hal.allocator.compute_size %allocator_8, shape = [], element_type = 16777248
    %allocator_10 = hal.buffer.allocator %1 : !hal.allocator
    %sz_11 = hal.allocator.compute_size %allocator_10, shape = [], element_type = 16777248
    %allocator_12 = hal.buffer.allocator %buffer : !hal.allocator
    %sz_13 = hal.allocator.compute_size %allocator_12, shape = [], element_type = 16777224
    hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%2, %c0, %sz_9), 1 = (%1, %c0, %sz_11), 2 = (%buffer, %c0, %sz_13)]
    hal.device.switch(%dev_6 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %exe : !hal.executable) {
      %c1_29 = constant 1 : index
      %c1_30 = constant 1 : index
      %c1_31 = constant 1 : index
      %c1_32 = constant 1 : index
      %7 = subi %c1_29, %c1_32 : index
      %8 = addi %c1_32, %7 : index
      %9 = divi_signed %8, %c1_29 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%9, %c1_32, %c1_32]
      hal.return
    }
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev_6, %cmd
    %allocator_14 = hal.buffer.allocator %buffer : !hal.allocator
    %off = hal.allocator.compute_offset %allocator_14, shape = [], element_type = 16777224, indices = []
    %3 = hal.buffer.load %buffer[%off] : i8
    %4 = trunci %3 : i8 to i1
    cond_br %4, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
  ^bb2(%5: !hal.buffer):  // pred: ^bb1
    %dev_15 = hal.ex.shared_device : !hal.device
    %allocator_16 = hal.device.allocator %dev_15 : !hal.allocator
    %sz_17 = hal.allocator.compute_size %allocator_16, shape = [], element_type = 16777248
    %buffer_18 = hal.allocator.allocate %allocator_16, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_17 : !hal.buffer
    hal.ex.defer_release %buffer_18 : !hal.buffer
    %cmd_19 = hal.command_buffer.create %dev_15, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_19
    %exe_20 = hal.executable.lookup %dev_15, @while_ex_dispatch_1 : !hal.executable
    %executable_layout_21 = hal.executable_layout.lookup %dev_15, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
    %c0_22 = constant 0 : index
    %allocator_23 = hal.buffer.allocator %5 : !hal.allocator
    %sz_24 = hal.allocator.compute_size %allocator_23, shape = [], element_type = 16777248
    %allocator_25 = hal.buffer.allocator %buffer_18 : !hal.allocator
    %sz_26 = hal.allocator.compute_size %allocator_25, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_19, %executable_layout_21, set=0, bindings=[0 = (%5, %c0_22, %sz_24), 1 = (%buffer_18, %c0_22, %sz_26)]
    hal.device.switch(%dev_15 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_19 : !hal.command_buffer, %arg2 = %exe_20 : !hal.executable) {
      %c1_29 = constant 1 : index
      %c1_30 = constant 1 : index
      %c1_31 = constant 1 : index
      %c1_32 = constant 1 : index
      %7 = subi %c1_29, %c1_32 : index
      %8 = addi %c1_32, %7 : index
      %9 = divi_signed %8, %c1_29 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%9, %c1_32, %c1_32]
      hal.return
    }
    %memory_barrier_27 = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd_19, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier_27]
    hal.command_buffer.end %cmd_19
    hal.ex.submit_and_wait %dev_15, %cmd_19
    br ^bb1(%buffer_18 : !hal.buffer)
  ^bb3(%6: !hal.buffer):  // pred: ^bb1
    %view = hal.buffer_view.create %6, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_28 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_28) : !hal.buffer_view
    return
  }
}

// *** IR Dump After mlir::iree_compiler::Shape::(anonymous namespace)::ExpandFunctionRankedShapeDimsPass ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %dev_0 = hal.ex.shared_device : !hal.device
  %allocator_1 = hal.device.allocator %dev_0 : !hal.allocator
  %cbuffer_2 = hal.allocator.allocate.const %allocator_1, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_2 : !hal.buffer
  %dev_3 = hal.ex.shared_device : !hal.device
  %allocator_4 = hal.device.allocator %dev_3 : !hal.allocator
  %cbuffer_5 = hal.allocator.allocate.const %allocator_4, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_5 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_5) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_2) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
  %dev_6 = hal.ex.shared_device : !hal.device
  %allocator_7 = hal.device.allocator %dev_6 : !hal.allocator
  %sz = hal.allocator.compute_size %allocator_7, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator_7, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev_6, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %exe = hal.executable.lookup %dev_6, @while_ex_dispatch_0 : !hal.executable
  %executable_layout = hal.executable_layout.lookup %dev_6, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %c0 = constant 0 : index
  %allocator_8 = hal.buffer.allocator %2 : !hal.allocator
  %sz_9 = hal.allocator.compute_size %allocator_8, shape = [], element_type = 16777248
  %allocator_10 = hal.buffer.allocator %1 : !hal.allocator
  %sz_11 = hal.allocator.compute_size %allocator_10, shape = [], element_type = 16777248
  %allocator_12 = hal.buffer.allocator %buffer : !hal.allocator
  %sz_13 = hal.allocator.compute_size %allocator_12, shape = [], element_type = 16777224
  hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%2, %c0, %sz_9), 1 = (%1, %c0, %sz_11), 2 = (%buffer, %c0, %sz_13)]
  hal.device.switch(%dev_6 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %exe : !hal.executable) {
    %c1_29 = constant 1 : index
    %c1_30 = constant 1 : index
    %c1_31 = constant 1 : index
    %c1_32 = constant 1 : index
    %7 = subi %c1_29, %c1_32 : index
    %8 = addi %c1_32, %7 : index
    %9 = divi_signed %8, %c1_29 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%9, %c1_32, %c1_32]
    hal.return
  }
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev_6, %cmd
  %allocator_14 = hal.buffer.allocator %buffer : !hal.allocator
  %off = hal.allocator.compute_offset %allocator_14, shape = [], element_type = 16777224, indices = []
  %3 = hal.buffer.load %buffer[%off] : i8
  %4 = trunci %3 : i8 to i1
  cond_br %4, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
^bb2(%5: !hal.buffer):  // pred: ^bb1
  %dev_15 = hal.ex.shared_device : !hal.device
  %allocator_16 = hal.device.allocator %dev_15 : !hal.allocator
  %sz_17 = hal.allocator.compute_size %allocator_16, shape = [], element_type = 16777248
  %buffer_18 = hal.allocator.allocate %allocator_16, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_17 : !hal.buffer
  hal.ex.defer_release %buffer_18 : !hal.buffer
  %cmd_19 = hal.command_buffer.create %dev_15, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_19
  %exe_20 = hal.executable.lookup %dev_15, @while_ex_dispatch_1 : !hal.executable
  %executable_layout_21 = hal.executable_layout.lookup %dev_15, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %c0_22 = constant 0 : index
  %allocator_23 = hal.buffer.allocator %5 : !hal.allocator
  %sz_24 = hal.allocator.compute_size %allocator_23, shape = [], element_type = 16777248
  %allocator_25 = hal.buffer.allocator %buffer_18 : !hal.allocator
  %sz_26 = hal.allocator.compute_size %allocator_25, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_19, %executable_layout_21, set=0, bindings=[0 = (%5, %c0_22, %sz_24), 1 = (%buffer_18, %c0_22, %sz_26)]
  hal.device.switch(%dev_15 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_19 : !hal.command_buffer, %arg2 = %exe_20 : !hal.executable) {
    %c1_29 = constant 1 : index
    %c1_30 = constant 1 : index
    %c1_31 = constant 1 : index
    %c1_32 = constant 1 : index
    %7 = subi %c1_29, %c1_32 : index
    %8 = addi %c1_32, %7 : index
    %9 = divi_signed %8, %c1_29 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%9, %c1_32, %c1_32]
    hal.return
  }
  %memory_barrier_27 = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd_19, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier_27]
  hal.command_buffer.end %cmd_19
  hal.ex.submit_and_wait %dev_15, %cmd_19
  br ^bb1(%buffer_18 : !hal.buffer)
^bb3(%6: !hal.buffer):  // pred: ^bb1
  %view = hal.buffer_view.create %6, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_28 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_28) : !hal.buffer_view
  return
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %c0 = constant 0 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %dev_0 = hal.ex.shared_device : !hal.device
  %allocator_1 = hal.device.allocator %dev_0 : !hal.allocator
  %cbuffer_2 = hal.allocator.allocate.const %allocator_1, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_2 : !hal.buffer
  %dev_3 = hal.ex.shared_device : !hal.device
  %allocator_4 = hal.device.allocator %dev_3 : !hal.allocator
  %cbuffer_5 = hal.allocator.allocate.const %allocator_4, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_5 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_5) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_2) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
  %dev_6 = hal.ex.shared_device : !hal.device
  %allocator_7 = hal.device.allocator %dev_6 : !hal.allocator
  %sz = hal.allocator.compute_size %allocator_7, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator_7, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev_6, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %exe = hal.executable.lookup %dev_6, @while_ex_dispatch_0 : !hal.executable
  %executable_layout = hal.executable_layout.lookup %dev_6, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %allocator_8 = hal.buffer.allocator %2 : !hal.allocator
  %sz_9 = hal.allocator.compute_size %allocator_8, shape = [], element_type = 16777248
  %allocator_10 = hal.buffer.allocator %1 : !hal.allocator
  %sz_11 = hal.allocator.compute_size %allocator_10, shape = [], element_type = 16777248
  %sz_12 = hal.allocator.compute_size %allocator_7, shape = [], element_type = 16777224
  hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%2, %c0, %sz_9), 1 = (%1, %c0, %sz_11), 2 = (%buffer, %c0, %sz_12)]
  hal.device.switch(%dev_6 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %exe : !hal.executable) {
    %c1_25 = constant 1 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_25, %c1_25, %c1_25]
    hal.return
  }
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev_6, %cmd
  %off = hal.allocator.compute_offset %allocator_7, shape = [], element_type = 16777224, indices = []
  %3 = hal.buffer.load %buffer[%off] : i8
  %4 = trunci %3 : i8 to i1
  cond_br %4, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
^bb2(%5: !hal.buffer):  // pred: ^bb1
  %dev_13 = hal.ex.shared_device : !hal.device
  %allocator_14 = hal.device.allocator %dev_13 : !hal.allocator
  %sz_15 = hal.allocator.compute_size %allocator_14, shape = [], element_type = 16777248
  %buffer_16 = hal.allocator.allocate %allocator_14, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_15 : !hal.buffer
  hal.ex.defer_release %buffer_16 : !hal.buffer
  %cmd_17 = hal.command_buffer.create %dev_13, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_17
  %exe_18 = hal.executable.lookup %dev_13, @while_ex_dispatch_1 : !hal.executable
  %executable_layout_19 = hal.executable_layout.lookup %dev_13, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %allocator_20 = hal.buffer.allocator %5 : !hal.allocator
  %sz_21 = hal.allocator.compute_size %allocator_20, shape = [], element_type = 16777248
  %sz_22 = hal.allocator.compute_size %allocator_14, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_17, %executable_layout_19, set=0, bindings=[0 = (%5, %c0, %sz_21), 1 = (%buffer_16, %c0, %sz_22)]
  hal.device.switch(%dev_13 : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_17 : !hal.command_buffer, %arg2 = %exe_18 : !hal.executable) {
    %c1_25 = constant 1 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_25, %c1_25, %c1_25]
    hal.return
  }
  %memory_barrier_23 = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd_17, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier_23]
  hal.command_buffer.end %cmd_17
  hal.ex.submit_and_wait %dev_13, %cmd_17
  br ^bb1(%buffer_16 : !hal.buffer)
^bb3(%6: !hal.buffer):  // pred: ^bb1
  %view = hal.buffer_view.create %6, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_24 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_24) : !hal.buffer_view
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export, iree.reflection = {f = "I1!R1!", fv = "1"}} {
  %c1 = constant 1 : index
  %c0 = constant 0 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_0 : !hal.buffer
  %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_1 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
  %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %exe = hal.executable.lookup %dev, @while_ex_dispatch_0 : !hal.executable
  %executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
  %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
  %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
  %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
  hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %exe : !hal.executable) {
    %c1_14 = constant 1 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_14, %c1_14, %c1_14]
    hal.return
  }
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev, %cmd
  %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
  %3 = hal.buffer.load %buffer[%off] : i8
  %4 = trunci %3 : i8 to i1
  cond_br %4, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
^bb2(%5: !hal.buffer):  // pred: ^bb1
  %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
  %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
  hal.ex.defer_release %buffer_7 : !hal.buffer
  %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_8
  %exe_9 = hal.executable.lookup %dev, @while_ex_dispatch_1 : !hal.executable
  %executable_layout_10 = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
  %allocator_11 = hal.buffer.allocator %5 : !hal.allocator
  %sz_12 = hal.allocator.compute_size %allocator_11, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_8, %executable_layout_10, set=0, bindings=[0 = (%5, %c0, %sz_12), 1 = (%buffer_7, %c0, %sz_6)]
  hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_8 : !hal.command_buffer, %arg2 = %exe_9 : !hal.executable) {
    %c1_14 = constant 1 : index
    hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_14, %c1_14, %c1_14]
    hal.return
  }
  hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd_8
  hal.ex.submit_and_wait %dev, %cmd_8
  br ^bb1(%buffer_7 : !hal.buffer)
^bb3(%6: !hal.buffer):  // pred: ^bb1
  %view = hal.buffer_view.create %6, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_13 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_13) : !hal.buffer_view
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::(anonymous namespace)::PublicABIGenerationPass ***


module {
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 1 : i32
            %1 = spv.constant 0 : i32
            %2 = spv.constant 8 : i32
            %3 = spv.constant 255 : i32
            %4 = spv.constant 4 : i32
            %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %9 = spv.Load "StorageBuffer" %8 : i32
            %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %11 = spv.Load "StorageBuffer" %10 : i32
            %12 = spv.SLessThan %9, %11 : i32
            %13 = spv.Select %12, %0, %1 : i1, i32
            %14 = spv.SMod %1, %4 : i32
            %15 = spv.IMul %14, %2 : i32
            %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
            %17 = spv.Not %16 : i32
            %18 = spv.BitwiseAnd %13, %3 : i32
            %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
            %20 = spv.SDiv %1, %4 : i32
            %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
            %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_0
          spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 0 : i32
            %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %4 = spv.Load "StorageBuffer" %3 : i32
            %5 = spv.Load "StorageBuffer" %3 : i32
            %6 = spv.IAdd %4, %5 : i32
            %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            spv.Store "StorageBuffer" %7, %6 : i32
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_1
          spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export = "while$raw"} {
    %c1 = constant 1 : index
    %c0 = constant 0 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_0 : !hal.buffer
    %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_1 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
    %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %exe = hal.executable.lookup %dev, @while_ex_dispatch_0 : !hal.executable
    %executable_layout = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
    %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
    %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
    %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
    %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd, %executable_layout, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
    hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %exe : !hal.executable) {
      %c1_14 = constant 1 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_14, %c1_14, %c1_14]
      hal.return
    }
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev, %cmd
    %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
    %3 = hal.buffer.load %buffer[%off] : i8
    %4 = trunci %3 : i8 to i1
    cond_br %4, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
  ^bb2(%5: !hal.buffer):  // pred: ^bb1
    %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
    %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
    hal.ex.defer_release %buffer_7 : !hal.buffer
    %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_8
    %exe_9 = hal.executable.lookup %dev, @while_ex_dispatch_1 : !hal.executable
    %executable_layout_10 = hal.executable_layout.lookup %dev, set_layouts = [[#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">]] : !hal.executable_layout
    %allocator_11 = hal.buffer.allocator %5 : !hal.allocator
    %sz_12 = hal.allocator.compute_size %allocator_11, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_8, %executable_layout_10, set=0, bindings=[0 = (%5, %c0, %sz_12), 1 = (%buffer_7, %c0, %sz_6)]
    hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_8 : !hal.command_buffer, %arg2 = %exe_9 : !hal.executable) {
      %c1_14 = constant 1 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_14, %c1_14, %c1_14]
      hal.return
    }
    hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd_8
    hal.ex.submit_and_wait %dev, %cmd_8
    br ^bb1(%buffer_7 : !hal.buffer)
  ^bb3(%6: !hal.buffer):  // pred: ^bb1
    %view = hal.buffer_view.create %6, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_13 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_13) : !hal.buffer_view
    return
  }
  func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
    call @while() : () -> ()
    return
  }
  func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
    %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
    hal.check_success %0, "semaphore wait failed"
    call @while$sync() : () -> ()
    hal.semaphore.signal %arg2, value = %arg3
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::MaterializeResourceCachesPass ***


module {
  hal.variable @_executable_while_ex_dispatch_0 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_1 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_descriptor_set_layout_1 init(@_descriptor_set_layout_1_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_1 init(@_executable_layout_1_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
  func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
    %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
    %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
    return %executable_cache_default : !hal.executable_cache
  }
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 1 : i32
            %1 = spv.constant 0 : i32
            %2 = spv.constant 8 : i32
            %3 = spv.constant 255 : i32
            %4 = spv.constant 4 : i32
            %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %9 = spv.Load "StorageBuffer" %8 : i32
            %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %11 = spv.Load "StorageBuffer" %10 : i32
            %12 = spv.SLessThan %9, %11 : i32
            %13 = spv.Select %12, %0, %1 : i1, i32
            %14 = spv.SMod %1, %4 : i32
            %15 = spv.IMul %14, %2 : i32
            %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
            %17 = spv.Not %16 : i32
            %18 = spv.BitwiseAnd %13, %3 : i32
            %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
            %20 = spv.SDiv %1, %4 : i32
            %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
            %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_0
          spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 0 : i32
            %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %4 = spv.Load "StorageBuffer" %3 : i32
            %5 = spv.Load "StorageBuffer" %3 : i32
            %6 = spv.IAdd %4, %5 : i32
            %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            spv.Store "StorageBuffer" %7, %6 : i32
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_1
          spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export = "while$raw"} {
    %c1 = constant 1 : index
    %c0 = constant 0 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_0 : !hal.buffer
    %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_1 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb2
    %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
    %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
    %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
    %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
    %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
    hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd : !hal.command_buffer, %arg2 = %3 : !hal.executable) {
      %c1_12 = constant 1 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_12, %c1_12, %c1_12]
      hal.return
    }
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev, %cmd
    %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
    %5 = hal.buffer.load %buffer[%off] : i8
    %6 = trunci %5 : i8 to i1
    cond_br %6, ^bb2(%2 : !hal.buffer), ^bb3(%2 : !hal.buffer)
  ^bb2(%7: !hal.buffer):  // pred: ^bb1
    %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
    %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
    hal.ex.defer_release %buffer_7 : !hal.buffer
    %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_8
    %8 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
    %9 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %allocator_9 = hal.buffer.allocator %7 : !hal.allocator
    %sz_10 = hal.allocator.compute_size %allocator_9, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_8, %9, set=0, bindings=[0 = (%7, %c0, %sz_10), 1 = (%buffer_7, %c0, %sz_6)]
    hal.device.switch(%dev : !hal.device)
    #hal.device.match.id<"vulkan*">(%arg0 = %c1 : index, %arg1 = %cmd_8 : !hal.command_buffer, %arg2 = %8 : !hal.executable) {
      %c1_12 = constant 1 : index
      hal.command_buffer.dispatch %arg1, %arg2, entry_point = 0, workgroup_xyz = [%c1_12, %c1_12, %c1_12]
      hal.return
    }
    hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd_8
    hal.ex.submit_and_wait %dev, %cmd_8
    br ^bb1(%buffer_7 : !hal.buffer)
  ^bb3(%10: !hal.buffer):  // pred: ^bb1
    %view = hal.buffer_view.create %10, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_11 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_11) : !hal.buffer_view
    return
  }
  func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
    call @while() : () -> ()
    return
  }
  func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
    %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
    hal.check_success %0, "semaphore wait failed"
    call @while$sync() : () -> ()
    hal.semaphore.signal %arg2, value = %arg3
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
  %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
  %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
  return %executable_cache_default : !hal.executable_cache
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @while() attributes {iree.module.export = "while$raw"} {
  %c1 = constant 1 : index
  %c0 = constant 0 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_0 : !hal.buffer
  %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_1 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb8
  %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
  %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
  %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
  %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
  %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
  %5 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
  cond_br %5, ^bb2, ^bb3
^bb2:  // pred: ^bb1
  %c1_6 = constant 1 : index
  hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1_6, %c1_6, %c1_6]
  br ^bb4
^bb3:  // pred: ^bb1
  iree.unreachable
^bb4:  // pred: ^bb2
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev, %cmd
  %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
  %6 = hal.buffer.load %buffer[%off] : i8
  %7 = trunci %6 : i8 to i1
  cond_br %7, ^bb5(%2 : !hal.buffer), ^bb9(%2 : !hal.buffer)
^bb5(%8: !hal.buffer):  // pred: ^bb4
  %sz_7 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
  %buffer_8 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_7 : !hal.buffer
  hal.ex.defer_release %buffer_8 : !hal.buffer
  %cmd_9 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_9
  %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
  %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %allocator_10 = hal.buffer.allocator %8 : !hal.allocator
  %sz_11 = hal.allocator.compute_size %allocator_10, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_9, %10, set=0, bindings=[0 = (%8, %c0, %sz_11), 1 = (%buffer_8, %c0, %sz_7)]
  %11 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
  cond_br %11, ^bb6, ^bb7
^bb6:  // pred: ^bb5
  %c1_12 = constant 1 : index
  hal.command_buffer.dispatch %cmd_9, %9, entry_point = 0, workgroup_xyz = [%c1_12, %c1_12, %c1_12]
  br ^bb8
^bb7:  // pred: ^bb5
  iree.unreachable
^bb8:  // pred: ^bb6
  hal.command_buffer.execution_barrier %cmd_9, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd_9
  hal.ex.submit_and_wait %dev, %cmd_9
  br ^bb1(%buffer_8 : !hal.buffer)
^bb9(%12: !hal.buffer):  // pred: ^bb4
  %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_13 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_13) : !hal.buffer_view
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
  call @while() : () -> ()
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::InlineDeviceSwitchesPass ***
func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
  %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
  hal.check_success %0, "semaphore wait failed"
  call @while$sync() : () -> ()
  hal.semaphore.signal %arg2, value = %arg3
  return
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::MemoizeDeviceQueriesPass ***


module {
  func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %0 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
    return %0 : i1
  }
  hal.variable @_device_match_id_0 init(@_device_match_id_0_initializer) : i1 attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_0 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_1 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_descriptor_set_layout_1 init(@_descriptor_set_layout_1_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_1 init(@_executable_layout_1_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
  func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
    %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
    %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
    return %executable_cache_default : !hal.executable_cache
  }
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_2__ bind(0, 2) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_0() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 1 : i32
            %1 = spv.constant 0 : i32
            %2 = spv.constant 8 : i32
            %3 = spv.constant 255 : i32
            %4 = spv.constant 4 : i32
            %5 = spv._address_of @__resource_var_0_2__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %6 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %7 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %8 = spv.AccessChain %6[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %9 = spv.Load "StorageBuffer" %8 : i32
            %10 = spv.AccessChain %7[%1, %1] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %11 = spv.Load "StorageBuffer" %10 : i32
            %12 = spv.SLessThan %9, %11 : i32
            %13 = spv.Select %12, %0, %1 : i1, i32
            %14 = spv.SMod %1, %4 : i32
            %15 = spv.IMul %14, %2 : i32
            %16 = spv.ShiftLeftLogical %3, %15 : i32, i32
            %17 = spv.Not %16 : i32
            %18 = spv.BitwiseAnd %13, %3 : i32
            %19 = spv.ShiftLeftLogical %18, %15 : i32, i32
            %20 = spv.SDiv %1, %4 : i32
            %21 = spv.AccessChain %5[%1, %20] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %22 = spv.AtomicAnd "Device" "AcquireRelease" %21, %17 : !spv.ptr<i32, StorageBuffer>
            %23 = spv.AtomicOr "Device" "AcquireRelease" %21, %19 : !spv.ptr<i32, StorageBuffer>
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_0
          spv.ExecutionMode @while_ex_dispatch_0 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.target "vulkan*" {
      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>}>} {
        spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
          spv.globalVariable @__resource_var_0_0__ bind(0, 0) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.globalVariable @__resource_var_0_1__ bind(0, 1) : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
          spv.func @while_ex_dispatch_1() "None" attributes {vkspv.workgroup_count_from_result_shape = 1 : i32} {
            %0 = spv.constant 0 : i32
            %1 = spv._address_of @__resource_var_0_1__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %2 = spv._address_of @__resource_var_0_0__ : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>
            %3 = spv.AccessChain %2[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            %4 = spv.Load "StorageBuffer" %3 : i32
            %5 = spv.Load "StorageBuffer" %3 : i32
            %6 = spv.IAdd %4, %5 : i32
            %7 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<!spv.array<1 x i32, stride=4> [0]>, StorageBuffer>, i32, i32
            spv.Store "StorageBuffer" %7, %6 : i32
            spv.Return
          }
          spv.EntryPoint "GLCompute" @while_ex_dispatch_1
          spv.ExecutionMode @while_ex_dispatch_1 "LocalSize", 1, 1, 1
        }
        hal.interface @legacy_io attributes {sym_visibility = "private"} {
          hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
          hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
        }
      }
    }
  }
  func @while() attributes {iree.module.export = "while$raw"} {
    %c1 = constant 1 : index
    %c0 = constant 0 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_0 : !hal.buffer
    %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_1 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb8
    %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
    %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
    %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
    %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
    %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
    %5 = hal.variable.load @_device_match_id_0 : i1
    cond_br %5, ^bb2, ^bb3
  ^bb2:  // pred: ^bb1
    %c1_6 = constant 1 : index
    hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1_6, %c1_6, %c1_6]
    br ^bb4
  ^bb3:  // pred: ^bb1
    iree.unreachable
  ^bb4:  // pred: ^bb2
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev, %cmd
    %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
    %6 = hal.buffer.load %buffer[%off] : i8
    %7 = trunci %6 : i8 to i1
    cond_br %7, ^bb5(%2 : !hal.buffer), ^bb9(%2 : !hal.buffer)
  ^bb5(%8: !hal.buffer):  // pred: ^bb4
    %sz_7 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
    %buffer_8 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_7 : !hal.buffer
    hal.ex.defer_release %buffer_8 : !hal.buffer
    %cmd_9 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_9
    %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
    %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %allocator_10 = hal.buffer.allocator %8 : !hal.allocator
    %sz_11 = hal.allocator.compute_size %allocator_10, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_9, %10, set=0, bindings=[0 = (%8, %c0, %sz_11), 1 = (%buffer_8, %c0, %sz_7)]
    %11 = hal.variable.load @_device_match_id_0 : i1
    cond_br %11, ^bb6, ^bb7
  ^bb6:  // pred: ^bb5
    %c1_12 = constant 1 : index
    hal.command_buffer.dispatch %cmd_9, %9, entry_point = 0, workgroup_xyz = [%c1_12, %c1_12, %c1_12]
    br ^bb8
  ^bb7:  // pred: ^bb5
    iree.unreachable
  ^bb8:  // pred: ^bb6
    hal.command_buffer.execution_barrier %cmd_9, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd_9
    hal.ex.submit_and_wait %dev, %cmd_9
    br ^bb1(%buffer_8 : !hal.buffer)
  ^bb9(%12: !hal.buffer):  // pred: ^bb4
    %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_13 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_13) : !hal.buffer_view
    return
  }
  func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
    call @while() : () -> ()
    return
  }
  func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
    %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
    hal.check_success %0, "semaphore wait failed"
    call @while$sync() : () -> ()
    hal.semaphore.signal %arg2, value = %arg3
    return
  }
}

// *** IR Dump After Canonicalizer ***
func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %0 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
  return %0 : i1
}

// *** IR Dump After CSE ***
func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %0 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
  return %0 : i1
}

// *** IR Dump After Canonicalizer ***
func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After CSE ***
func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After Canonicalizer ***
func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After CSE ***
func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After Canonicalizer ***
func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After CSE ***
func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
  return %descriptor_set_layout : !hal.descriptor_set_layout
}

// *** IR Dump After Canonicalizer ***
func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After CSE ***
func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
  %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
  %dev = hal.ex.shared_device : !hal.device
  %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
  return %executable_layout : !hal.executable_layout
}

// *** IR Dump After Canonicalizer ***
func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
  %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
  %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
  return %executable_cache_default : !hal.executable_cache
}

// *** IR Dump After CSE ***
func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
  %dev = hal.ex.shared_device : !hal.device
  %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
  %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
  %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
  hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
  return %executable_cache_default : !hal.executable_cache
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::SerializeExecutablesPass ***
hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
  hal.interface @legacy_io {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
  }
  hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
  hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>, format = 1397773893 : i32} {
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::HAL::SerializeExecutablesPass ***
hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
  hal.interface @legacy_io {
    hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
    hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
  }
  hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
  hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>, format = 1397773893 : i32} {
  }
}

// *** IR Dump After Canonicalizer ***
func @while() attributes {iree.module.export = "while$raw"} {
  %c0 = constant 0 : index
  %c1 = constant 1 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_0 : !hal.buffer
  %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_1 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb5
  %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
  %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
  %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
  %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
  %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
  %5 = hal.variable.load @_device_match_id_0 : i1
  cond_br %5, ^bb2, ^bb3
^bb2:  // pred: ^bb1
  hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev, %cmd
  %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
  %6 = hal.buffer.load %buffer[%off] : i8
  %7 = trunci %6 : i8 to i1
  cond_br %7, ^bb4(%2 : !hal.buffer), ^bb6(%2 : !hal.buffer)
^bb3:  // 2 preds: ^bb1, ^bb4
  iree.unreachable
^bb4(%8: !hal.buffer):  // pred: ^bb2
  %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
  %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
  hal.ex.defer_release %buffer_7 : !hal.buffer
  %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_8
  %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
  %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %allocator_9 = hal.buffer.allocator %8 : !hal.allocator
  %sz_10 = hal.allocator.compute_size %allocator_9, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_8, %10, set=0, bindings=[0 = (%8, %c0, %sz_10), 1 = (%buffer_7, %c0, %sz_6)]
  %11 = hal.variable.load @_device_match_id_0 : i1
  cond_br %11, ^bb5, ^bb3
^bb5:  // pred: ^bb4
  hal.command_buffer.dispatch %cmd_8, %9, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
  hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd_8
  hal.ex.submit_and_wait %dev, %cmd_8
  br ^bb1(%buffer_7 : !hal.buffer)
^bb6(%12: !hal.buffer):  // pred: ^bb2
  %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_11 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_11) : !hal.buffer_view
  return
}

// *** IR Dump After CSE ***
func @while() attributes {iree.module.export = "while$raw"} {
  %c0 = constant 0 : index
  %c1 = constant 1 : index
  %dev = hal.ex.shared_device : !hal.device
  %allocator = hal.device.allocator %dev : !hal.allocator
  %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
  hal.ex.defer_release %cbuffer : !hal.buffer
  %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
  hal.ex.defer_release %cbuffer_0 : !hal.buffer
  %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
  hal.ex.defer_release %cbuffer_1 : !hal.buffer
  %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
  %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
  br ^bb1(%0 : !hal.buffer)
^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb5
  %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
  %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
  hal.ex.defer_release %buffer : !hal.buffer
  %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd
  %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
  %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
  %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
  %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
  %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
  %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
  %5 = hal.variable.load @_device_match_id_0 : i1
  cond_br %5, ^bb2, ^bb3
^bb2:  // pred: ^bb1
  hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
  %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
  hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd
  hal.ex.submit_and_wait %dev, %cmd
  %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
  %6 = hal.buffer.load %buffer[%off] : i8
  %7 = trunci %6 : i8 to i1
  cond_br %7, ^bb4(%2 : !hal.buffer), ^bb6(%2 : !hal.buffer)
^bb3:  // 2 preds: ^bb1, ^bb4
  iree.unreachable
^bb4(%8: !hal.buffer):  // pred: ^bb2
  %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
  %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
  hal.ex.defer_release %buffer_7 : !hal.buffer
  %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
  hal.command_buffer.begin %cmd_8
  %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
  %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
  %allocator_9 = hal.buffer.allocator %8 : !hal.allocator
  %sz_10 = hal.allocator.compute_size %allocator_9, shape = [], element_type = 16777248
  hal.command_buffer.push_descriptor_set %cmd_8, %10, set=0, bindings=[0 = (%8, %c0, %sz_10), 1 = (%buffer_7, %c0, %sz_6)]
  %11 = hal.variable.load @_device_match_id_0 : i1
  cond_br %11, ^bb5, ^bb3
^bb5:  // pred: ^bb4
  hal.command_buffer.dispatch %cmd_8, %9, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
  hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
  hal.command_buffer.end %cmd_8
  hal.ex.submit_and_wait %dev, %cmd_8
  br ^bb1(%buffer_7 : !hal.buffer)
^bb6(%12: !hal.buffer):  // pred: ^bb2
  %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
  %view_11 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
  check.expect_eq(%view, %view_11) : !hal.buffer_view
  return
}

// *** IR Dump After Canonicalizer ***
func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
  call @while() : () -> ()
  return
}

// *** IR Dump After CSE ***
func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
  call @while() : () -> ()
  return
}

// *** IR Dump After Canonicalizer ***
func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
  %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
  hal.check_success %0, "semaphore wait failed"
  call @while$sync() : () -> ()
  hal.semaphore.signal %arg2, value = %arg3
  return
}

// *** IR Dump After CSE ***
func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
  %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
  hal.check_success %0, "semaphore wait failed"
  call @while$sync() : () -> ()
  hal.semaphore.signal %arg2, value = %arg3
  return
}

// *** IR Dump After SymbolDCE ***


module {
  func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %0 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
    return %0 : i1
  }
  hal.variable @_device_match_id_0 init(@_device_match_id_0_initializer) : i1 attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_0 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_1 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_descriptor_set_layout_1 init(@_descriptor_set_layout_1_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_1 init(@_executable_layout_1_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
  func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
    %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
    %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
    return %executable_cache_default : !hal.executable_cache
  }
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>, format = 1397773893 : i32} {
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>, format = 1397773893 : i32} {
    }
  }
  func @while() attributes {iree.module.export = "while$raw"} {
    %c0 = constant 0 : index
    %c1 = constant 1 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_0 : !hal.buffer
    %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_1 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb5
    %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
    %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
    %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
    %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
    %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
    %5 = hal.variable.load @_device_match_id_0 : i1
    cond_br %5, ^bb2, ^bb3
  ^bb2:  // pred: ^bb1
    hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev, %cmd
    %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
    %6 = hal.buffer.load %buffer[%off] : i8
    %7 = trunci %6 : i8 to i1
    cond_br %7, ^bb4(%2 : !hal.buffer), ^bb6(%2 : !hal.buffer)
  ^bb3:  // 2 preds: ^bb1, ^bb4
    iree.unreachable
  ^bb4(%8: !hal.buffer):  // pred: ^bb2
    %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
    %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
    hal.ex.defer_release %buffer_7 : !hal.buffer
    %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_8
    %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
    %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %allocator_9 = hal.buffer.allocator %8 : !hal.allocator
    %sz_10 = hal.allocator.compute_size %allocator_9, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_8, %10, set=0, bindings=[0 = (%8, %c0, %sz_10), 1 = (%buffer_7, %c0, %sz_6)]
    %11 = hal.variable.load @_device_match_id_0 : i1
    cond_br %11, ^bb5, ^bb3
  ^bb5:  // pred: ^bb4
    hal.command_buffer.dispatch %cmd_8, %9, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
    hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd_8
    hal.ex.submit_and_wait %dev, %cmd_8
    br ^bb1(%buffer_7 : !hal.buffer)
  ^bb6(%12: !hal.buffer):  // pred: ^bb2
    %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_11 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_11) : !hal.buffer_view
    return
  }
  func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
    call @while() : () -> ()
    return
  }
  func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
    %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
    hal.check_success %0, "semaphore wait failed"
    call @while$sync() : () -> ()
    hal.semaphore.signal %arg2, value = %arg3
    return
  }
}

// *** IR Dump After Canonicalizer ***


module {
  func @_device_match_id_0_initializer() -> i1 attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %0 = hal.device.match.id %dev, pattern = ["vulkan*"] : (!hal.device) -> i1
    return %0 : i1
  }
  hal.variable @_device_match_id_0 init(@_device_match_id_0_initializer) : i1 attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_0 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_executable_while_ex_dispatch_1 mutable : !hal.executable attributes {sym_visibility = "private"}
  hal.variable @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_0_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<2, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_0 init(@_executable_layout_0_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_0_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_0 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_descriptor_set_layout_1 init(@_descriptor_set_layout_1_initializer) : !hal.descriptor_set_layout attributes {sym_visibility = "private"}
  func @_descriptor_set_layout_1_initializer() -> !hal.descriptor_set_layout attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %descriptor_set_layout = hal.descriptor_set_layout.create %dev, "PushOnly", bindings = [#hal.descriptor_set_layout_binding<0, "StorageBuffer", "Read">, #hal.descriptor_set_layout_binding<1, "StorageBuffer", "Write|Discard">] : !hal.descriptor_set_layout
    return %descriptor_set_layout : !hal.descriptor_set_layout
  }
  hal.variable @_executable_layout_1 init(@_executable_layout_1_initializer) : !hal.executable_layout attributes {sym_visibility = "private"}
  func @_executable_layout_1_initializer() -> !hal.executable_layout attributes {sym_visibility = "private"} {
    %0 = hal.variable.load @_descriptor_set_layout_1 : !hal.descriptor_set_layout
    %dev = hal.ex.shared_device : !hal.device
    %executable_layout = hal.executable_layout.create %dev, set_layouts = [%0], push_constants = 0 : !hal.executable_layout
    return %executable_layout : !hal.executable_layout
  }
  hal.variable @_executable_cache init(@_executable_cache_initializer) : !hal.executable_cache
  func @_executable_cache_initializer() -> !hal.executable_cache attributes {sym_visibility = "private"} {
    %dev = hal.ex.shared_device : !hal.device
    %executable_cache_default = hal.executable_cache.create %dev, identifier = "default" : !hal.executable_cache
    %0 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %executable_while_ex_dispatch_0 = hal.executable_cache.prepare %executable_cache_default, layout = %0, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_0 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_0, @_executable_while_ex_dispatch_0 : !hal.executable
    %1 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %executable_while_ex_dispatch_1 = hal.executable_cache.prepare %executable_cache_default, layout = %1, caching_mode = "AliasProvidedData|AllowPersistentCaching|AllowOptimization", @while_ex_dispatch_1 : !hal.executable
    hal.variable.store %executable_while_ex_dispatch_1, @_executable_while_ex_dispatch_1 : !hal.executable
    return %executable_cache_default : !hal.executable_cache
  }
  hal.executable @while_ex_dispatch_0 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @arg1, set=0, binding=1, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=2, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_0 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>, tensor<i32>) -> tensor<i8>}
    hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>, format = 1397773893 : i32} {
    }
  }
  hal.executable @while_ex_dispatch_1 attributes {sym_visibility = "private"} {
    hal.interface @legacy_io {
      hal.interface.binding @arg0, set=0, binding=0, type="StorageBuffer", access="Read"
      hal.interface.binding @ret0, set=0, binding=1, type="StorageBuffer", access="Write|Discard"
    }
    hal.executable.entry_point @while_ex_dispatch_1 attributes {interface = @legacy_io, ordinal = 0 : i32, signature = (tensor<i32>) -> tensor<i32>}
    hal.executable.binary attributes {data = dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>, format = 1397773893 : i32} {
    }
  }
  func @while() attributes {iree.module.export = "while$raw"} {
    %c0 = constant 0 : index
    %c1 = constant 1 : index
    %dev = hal.ex.shared_device : !hal.device
    %allocator = hal.device.allocator %dev : !hal.allocator
    %cbuffer = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<4> : tensor<i32>
    hal.ex.defer_release %cbuffer : !hal.buffer
    %cbuffer_0 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<3> : tensor<i32>
    hal.ex.defer_release %cbuffer_0 : !hal.buffer
    %cbuffer_1 = hal.allocator.allocate.const %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch" : !hal.buffer = dense<1> : tensor<i32>
    hal.ex.defer_release %cbuffer_1 : !hal.buffer
    %0 = iree.do_not_optimize(%cbuffer_1) : !hal.buffer
    %1 = iree.do_not_optimize(%cbuffer_0) : !hal.buffer
    br ^bb1(%0 : !hal.buffer)
  ^bb1(%2: !hal.buffer):  // 2 preds: ^bb0, ^bb5
    %sz = hal.allocator.compute_size %allocator, shape = [], element_type = 16777224
    %buffer = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz : !hal.buffer
    hal.ex.defer_release %buffer : !hal.buffer
    %cmd = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd
    %3 = hal.variable.load @_executable_while_ex_dispatch_0 : !hal.executable
    %4 = hal.variable.load @_executable_layout_0 : !hal.executable_layout
    %allocator_2 = hal.buffer.allocator %2 : !hal.allocator
    %sz_3 = hal.allocator.compute_size %allocator_2, shape = [], element_type = 16777248
    %allocator_4 = hal.buffer.allocator %1 : !hal.allocator
    %sz_5 = hal.allocator.compute_size %allocator_4, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd, %4, set=0, bindings=[0 = (%2, %c0, %sz_3), 1 = (%1, %c0, %sz_5), 2 = (%buffer, %c0, %sz)]
    %5 = hal.variable.load @_device_match_id_0 : i1
    cond_br %5, ^bb2, ^bb3
  ^bb2:  // pred: ^bb1
    hal.command_buffer.dispatch %cmd, %3, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
    %memory_barrier = hal.make_memory_barrier "DispatchWrite", "DispatchRead" : tuple<i32, i32>
    hal.command_buffer.execution_barrier %cmd, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd
    hal.ex.submit_and_wait %dev, %cmd
    %off = hal.allocator.compute_offset %allocator, shape = [], element_type = 16777224, indices = []
    %6 = hal.buffer.load %buffer[%off] : i8
    %7 = trunci %6 : i8 to i1
    cond_br %7, ^bb4(%2 : !hal.buffer), ^bb6(%2 : !hal.buffer)
  ^bb3:  // 2 preds: ^bb1, ^bb4
    iree.unreachable
  ^bb4(%8: !hal.buffer):  // pred: ^bb2
    %sz_6 = hal.allocator.compute_size %allocator, shape = [], element_type = 16777248
    %buffer_7 = hal.allocator.allocate %allocator, "HostVisible|DeviceVisible|DeviceLocal", "Constant|Transfer|Mapping|Dispatch", %sz_6 : !hal.buffer
    hal.ex.defer_release %buffer_7 : !hal.buffer
    %cmd_8 = hal.command_buffer.create %dev, "OneShot", "Transfer|Dispatch" : !hal.command_buffer
    hal.command_buffer.begin %cmd_8
    %9 = hal.variable.load @_executable_while_ex_dispatch_1 : !hal.executable
    %10 = hal.variable.load @_executable_layout_1 : !hal.executable_layout
    %allocator_9 = hal.buffer.allocator %8 : !hal.allocator
    %sz_10 = hal.allocator.compute_size %allocator_9, shape = [], element_type = 16777248
    hal.command_buffer.push_descriptor_set %cmd_8, %10, set=0, bindings=[0 = (%8, %c0, %sz_10), 1 = (%buffer_7, %c0, %sz_6)]
    %11 = hal.variable.load @_device_match_id_0 : i1
    cond_br %11, ^bb5, ^bb3
  ^bb5:  // pred: ^bb4
    hal.command_buffer.dispatch %cmd_8, %9, entry_point = 0, workgroup_xyz = [%c1, %c1, %c1]
    hal.command_buffer.execution_barrier %cmd_8, "Dispatch|CommandRetire", "CommandIssue|Dispatch", memory_barriers=[%memory_barrier]
    hal.command_buffer.end %cmd_8
    hal.ex.submit_and_wait %dev, %cmd_8
    br ^bb1(%buffer_7 : !hal.buffer)
  ^bb6(%12: !hal.buffer):  // pred: ^bb2
    %view = hal.buffer_view.create %12, shape = [], element_type = 16777248 : !hal.buffer_view
    %view_11 = hal.buffer_view.create %cbuffer, shape = [], element_type = 16777248 : !hal.buffer_view
    check.expect_eq(%view, %view_11) : !hal.buffer_view
    return
  }
  func @while$sync() attributes {iree.abi.stub, iree.module.export = "while", iree.reflection = {f = "I1!R1!", fv = "1"}} {
    call @while() : () -> ()
    return
  }
  func @while$async(%arg0: !hal.semaphore, %arg1: index, %arg2: !hal.semaphore, %arg3: index) {
    %0 = hal.semaphore.await %arg0, min_value = %arg1 : i32
    hal.check_success %0, "semaphore wait failed"
    call @while$sync() : () -> ()
    hal.semaphore.signal %arg2, value = %arg3
    return
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::VM::ConversionPass ***


module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.func @_device_match_id_0_initializer() -> i32 attributes {sym_visibility = "private"} {
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.return %0 : i32
    }
    vm.global.i32 @_device_match_id_0 init(@_device_match_id_0_initializer) : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 init(@_descriptor_set_layout_0_initializer) : !vm.ref<!hal.descriptor_set_layout>
    vm.func @_descriptor_set_layout_0_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} {
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c7 = vm.const.i32 7 : i32
      %c1_0 = vm.const.i32 1 : i32
      %c1_1 = vm.const.i32 1 : i32
      %c7_2 = vm.const.i32 7 : i32
      %c1_3 = vm.const.i32 1 : i32
      %c2 = vm.const.i32 2 : i32
      %c7_4 = vm.const.i32 7 : i32
      %c6 = vm.const.i32 6 : i32
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [(%zero, %c7, %c1_0), (%c1_1, %c7_2, %c1_3), (%c2, %c7_4, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.return %ref_5 : !vm.ref<!hal.descriptor_set_layout>
    }
    vm.global.ref @_executable_layout_0 init(@_executable_layout_0_initializer) : !vm.ref<!hal.executable_layout>
    vm.func @_executable_layout_0_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} {
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %zero = vm.const.i32.zero : i32
      %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.return %ref_0 : !vm.ref<!hal.executable_layout>
    }
    vm.global.ref @_descriptor_set_layout_1 init(@_descriptor_set_layout_1_initializer) : !vm.ref<!hal.descriptor_set_layout>
    vm.func @_descriptor_set_layout_1_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} {
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c7 = vm.const.i32 7 : i32
      %c1_0 = vm.const.i32 1 : i32
      %c1_1 = vm.const.i32 1 : i32
      %c7_2 = vm.const.i32 7 : i32
      %c6 = vm.const.i32 6 : i32
      %ref_3 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [(%zero, %c7, %c1_0), (%c1_1, %c7_2, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.return %ref_3 : !vm.ref<!hal.descriptor_set_layout>
    }
    vm.global.ref @_executable_layout_1 init(@_executable_layout_1_initializer) : !vm.ref<!hal.executable_layout>
    vm.func @_executable_layout_1_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} {
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %zero = vm.const.i32.zero : i32
      %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.return %ref_0 : !vm.ref<!hal.executable_layout>
    }
    vm.global.ref @_executable_cache init(@_executable_cache_initializer) : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.func @_executable_cache_initializer() -> !vm.ref<!hal.executable_cache> attributes {sym_visibility = "private"} {
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_0 = vm.call @hal.executable_cache.create(%ref, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %c1397773893 = vm.const.i32 1397773893 : i32
      %0 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.switch.ref %0[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref_2 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_0, %c7, %ref_1) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_2, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %c1397773893_3 = vm.const.i32 1397773893 : i32
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1397773893_3]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %null_4 = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %ref_5 = vm.switch.ref %1[%_while_ex_dispatch_1_binary_spirv] else %null_4 : !vm.ref<!iree.byte_buffer>
      %c7_6 = vm.const.i32 7 : i32
      %ref_7 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_1, %c7_6, %ref_5) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_7, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.return %ref_0 : !vm.ref<!hal.executable_cache>
    }
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %zero = vm.const.i32.zero : i32
      %c1 = vm.const.i32 1 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %c50_2 = vm.const.i32 50 : i32
      %c15_3 = vm.const.i32 15 : i32
      %c16777248_4 = vm.const.i32 16777248 : i32
      %ref_5 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50_2, %c15_3, [], %c16777248_4, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_5) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %c50_6 = vm.const.i32 50 : i32
      %c15_7 = vm.const.i32 15 : i32
      %c16777248_8 = vm.const.i32 16777248 : i32
      %ref_9 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50_6, %c15_7, [], %c16777248_8, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_9) : (!vm.ref<!hal.buffer>) -> ()
      %0 = iree.do_not_optimize(%ref_9) : !vm.ref<!hal.buffer>
      %1 = iree.do_not_optimize(%ref_5) : !vm.ref<!hal.buffer>
      vm.br ^bb1(%0 : !vm.ref<!hal.buffer>)
    ^bb1(%2: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %c16777224 = vm.const.i32 16777224 : i32
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %c50_10 = vm.const.i32 50 : i32
      %c15_11 = vm.const.i32 15 : i32
      %ref_12 = vm.call @hal.allocator.allocate(%ref_0, %c50_10, %c15_11, %3) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_12) : (!vm.ref<!hal.buffer>) -> ()
      %c1_13 = vm.const.i32 1 : i32
      %c3 = vm.const.i32 3 : i32
      %ref_14 = vm.call @hal.command_buffer.create(%ref, %c1_13, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_14) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_15 = vm.call @hal.buffer.allocator(%2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %c16777248_16 = vm.const.i32 16777248 : i32
      %4 = vm.call.variadic @hal.allocator.compute_size(%ref_15, [], %c16777248_16) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_17 = vm.call @hal.buffer.allocator(%1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %c16777248_18 = vm.const.i32 16777248 : i32
      %5 = vm.call.variadic @hal.allocator.compute_size(%ref_17, [], %c16777248_18) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %zero_19 = vm.const.i32.zero : i32
      %zero_20 = vm.const.i32.zero : i32
      %c1_21 = vm.const.i32 1 : i32
      %c2 = vm.const.i32 2 : i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_14, %_executable_layout_0, %zero_19, [%zero_20, %c1_21, %c2], [%2, %1, %ref_12], [%zero, %zero, %zero], [%4, %5, %3]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      %zero_22 = vm.const.i32.zero : i32
      vm.call @hal.command_buffer.dispatch(%ref_14, %_executable_while_ex_dispatch_0, %zero_22, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_14, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_14) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_14) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %c16777224_23 = vm.const.i32 16777224 : i32
      %6 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224_23, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %c1_24 = vm.const.i32 1 : i32
      %7 = vm.call @hal.buffer.load(%ref_12, %6, %c1_24) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %7, ^bb4(%2 : !vm.ref<!hal.buffer>), ^bb6(%2 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      %c2_25 = vm.const.i32 2 : i32
      vm.fail %c2_25, "unreachable location reached"
    ^bb4(%8: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %c16777248_26 = vm.const.i32 16777248 : i32
      %9 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248_26) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %c50_27 = vm.const.i32 50 : i32
      %c15_28 = vm.const.i32 15 : i32
      %ref_29 = vm.call @hal.allocator.allocate(%ref_0, %c50_27, %c15_28, %9) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_29) : (!vm.ref<!hal.buffer>) -> ()
      %c1_30 = vm.const.i32 1 : i32
      %c3_31 = vm.const.i32 3 : i32
      %ref_32 = vm.call @hal.command_buffer.create(%ref, %c1_30, %c3_31) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_32) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_33 = vm.call @hal.buffer.allocator(%8) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %c16777248_34 = vm.const.i32 16777248 : i32
      %10 = vm.call.variadic @hal.allocator.compute_size(%ref_33, [], %c16777248_34) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %zero_35 = vm.const.i32.zero : i32
      %zero_36 = vm.const.i32.zero : i32
      %c1_37 = vm.const.i32 1 : i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_32, %_executable_layout_1, %zero_35, [%zero_36, %c1_37], [%8, %ref_29], [%zero, %zero], [%10, %9]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_38 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_38, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      %zero_39 = vm.const.i32.zero : i32
      vm.call @hal.command_buffer.dispatch(%ref_32, %_executable_while_ex_dispatch_1, %zero_39, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      %c20_40 = vm.const.i32 20 : i32
      %c5_41 = vm.const.i32 5 : i32
      %c8_42 = vm.const.i32 8 : i32
      %c4_43 = vm.const.i32 4 : i32
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_32, %c20_40, %c5_41, [%c8_42], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_32) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_32) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_29 : !vm.ref<!hal.buffer>)
    ^bb6(%11: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %c16777248_44 = vm.const.i32 16777248 : i32
      %ref_45 = vm.call.variadic @hal.buffer_view.create(%11, [], %c16777248_44) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %c16777248_46 = vm.const.i32 16777248 : i32
      %ref_47 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248_46) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_45, %ref_47) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_fail %0, "semaphore wait failed"
      vm.call @while$sync() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    }
    vm.import @check.expect_true(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_false(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_all_true(%operand : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_almost_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.trace(%operands : !vm.ref<!hal.buffer_view>...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.create(%device : !vm.ref<!hal.device>, %initial_value : i32) -> !vm.ref<!hal.semaphore> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.query(%semaphore : !vm.ref<!hal.semaphore>) -> (i32, i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.fail(%semaphore : !vm.ref<!hal.semaphore>, %status : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::VM::GlobalInitializationPass ***
vm.module @module {
  vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
  vm.func @_device_match_id_0_initializer() -> i32 attributes {sym_visibility = "private"} {
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
    %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
    vm.return %0 : i32
  }
  vm.global.i32 @_device_match_id_0 mutable : i32
  vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
  vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
  vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
  vm.func @_descriptor_set_layout_0_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} {
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %c1 = vm.const.i32 1 : i32
    %zero = vm.const.i32.zero : i32
    %c7 = vm.const.i32 7 : i32
    %c1_0 = vm.const.i32 1 : i32
    %c1_1 = vm.const.i32 1 : i32
    %c7_2 = vm.const.i32 7 : i32
    %c1_3 = vm.const.i32 1 : i32
    %c2 = vm.const.i32 2 : i32
    %c7_4 = vm.const.i32 7 : i32
    %c6 = vm.const.i32 6 : i32
    %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [(%zero, %c7, %c1_0), (%c1_1, %c7_2, %c1_3), (%c2, %c7_4, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
    vm.return %ref_5 : !vm.ref<!hal.descriptor_set_layout>
  }
  vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
  vm.func @_executable_layout_0_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} {
    %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %zero = vm.const.i32.zero : i32
    %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
    vm.return %ref_0 : !vm.ref<!hal.executable_layout>
  }
  vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
  vm.func @_descriptor_set_layout_1_initializer() -> !vm.ref<!hal.descriptor_set_layout> attributes {sym_visibility = "private"} {
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %c1 = vm.const.i32 1 : i32
    %zero = vm.const.i32.zero : i32
    %c7 = vm.const.i32 7 : i32
    %c1_0 = vm.const.i32 1 : i32
    %c1_1 = vm.const.i32 1 : i32
    %c7_2 = vm.const.i32 7 : i32
    %c6 = vm.const.i32 6 : i32
    %ref_3 = vm.call.variadic @hal.descriptor_set_layout.create(%ref, %c1, [(%zero, %c7, %c1_0), (%c1_1, %c7_2, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
    vm.return %ref_3 : !vm.ref<!hal.descriptor_set_layout>
  }
  vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
  vm.func @_executable_layout_1_initializer() -> !vm.ref<!hal.executable_layout> attributes {sym_visibility = "private"} {
    %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %zero = vm.const.i32.zero : i32
    %ref_0 = vm.call.variadic @hal.executable_layout.create(%ref, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
    vm.return %ref_0 : !vm.ref<!hal.executable_layout>
  }
  vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
  vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
  vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
  vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
  vm.func @_executable_cache_initializer() -> !vm.ref<!hal.executable_cache> attributes {sym_visibility = "private"} {
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
    %ref_0 = vm.call @hal.executable_cache.create(%ref, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
    %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
    %c1397773893 = vm.const.i32 1397773893 : i32
    %0 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
    %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
    %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
    %ref_1 = vm.switch.ref %0[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
    %c7 = vm.const.i32 7 : i32
    %ref_2 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_0, %c7, %ref_1) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
    vm.global.store.ref %ref_2, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
    %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
    %c1397773893_3 = vm.const.i32 1397773893 : i32
    %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_0, [%c1397773893_3]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
    %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
    %null_4 = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
    %ref_5 = vm.switch.ref %1[%_while_ex_dispatch_1_binary_spirv] else %null_4 : !vm.ref<!iree.byte_buffer>
    %c7_6 = vm.const.i32 7 : i32
    %ref_7 = vm.call @hal.executable_cache.prepare(%ref_0, %_executable_layout_1, %c7_6, %ref_5) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
    vm.global.store.ref %ref_7, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
    vm.return %ref_0 : !vm.ref<!hal.executable_cache>
  }
  vm.rodata @while_const_0 dense<4> : tensor<i32>
  vm.rodata @while_const_1 dense<3> : tensor<i32>
  vm.rodata @while_const_2 dense<1> : tensor<i32>
  vm.func @while() {
    %zero = vm.const.i32.zero : i32
    %c1 = vm.const.i32 1 : i32
    %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
    %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
    %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
    %c50 = vm.const.i32 50 : i32
    %c15 = vm.const.i32 15 : i32
    %c16777248 = vm.const.i32 16777248 : i32
    %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
    vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
    %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
    %c50_2 = vm.const.i32 50 : i32
    %c15_3 = vm.const.i32 15 : i32
    %c16777248_4 = vm.const.i32 16777248 : i32
    %ref_5 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50_2, %c15_3, [], %c16777248_4, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
    vm.call @hal.ex.defer_release(%ref_5) : (!vm.ref<!hal.buffer>) -> ()
    %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
    %c50_6 = vm.const.i32 50 : i32
    %c15_7 = vm.const.i32 15 : i32
    %c16777248_8 = vm.const.i32 16777248 : i32
    %ref_9 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50_6, %c15_7, [], %c16777248_8, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
    vm.call @hal.ex.defer_release(%ref_9) : (!vm.ref<!hal.buffer>) -> ()
    %0 = iree.do_not_optimize(%ref_9) : !vm.ref<!hal.buffer>
    %1 = iree.do_not_optimize(%ref_5) : !vm.ref<!hal.buffer>
    vm.br ^bb1(%0 : !vm.ref<!hal.buffer>)
  ^bb1(%2: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
    %c16777224 = vm.const.i32 16777224 : i32
    %3 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
    %c50_10 = vm.const.i32 50 : i32
    %c15_11 = vm.const.i32 15 : i32
    %ref_12 = vm.call @hal.allocator.allocate(%ref_0, %c50_10, %c15_11, %3) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
    vm.call @hal.ex.defer_release(%ref_12) : (!vm.ref<!hal.buffer>) -> ()
    %c1_13 = vm.const.i32 1 : i32
    %c3 = vm.const.i32 3 : i32
    %ref_14 = vm.call @hal.command_buffer.create(%ref, %c1_13, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
    vm.call @hal.command_buffer.begin(%ref_14) : (!vm.ref<!hal.command_buffer>) -> ()
    %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
    %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
    %ref_15 = vm.call @hal.buffer.allocator(%2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
    %c16777248_16 = vm.const.i32 16777248 : i32
    %4 = vm.call.variadic @hal.allocator.compute_size(%ref_15, [], %c16777248_16) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
    %ref_17 = vm.call @hal.buffer.allocator(%1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
    %c16777248_18 = vm.const.i32 16777248 : i32
    %5 = vm.call.variadic @hal.allocator.compute_size(%ref_17, [], %c16777248_18) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
    %zero_19 = vm.const.i32.zero : i32
    %zero_20 = vm.const.i32.zero : i32
    %c1_21 = vm.const.i32 1 : i32
    %c2 = vm.const.i32 2 : i32
    vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_14, %_executable_layout_0, %zero_19, [%zero_20, %c1_21, %c2], [%2, %1, %ref_12], [%zero, %zero, %zero], [%4, %5, %3]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
    %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
    vm.cond_br %_device_match_id_0, ^bb2, ^bb3
  ^bb2:  // pred: ^bb1
    %zero_22 = vm.const.i32.zero : i32
    vm.call @hal.command_buffer.dispatch(%ref_14, %_executable_while_ex_dispatch_0, %zero_22, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
    %c20 = vm.const.i32 20 : i32
    %c5 = vm.const.i32 5 : i32
    %c8 = vm.const.i32 8 : i32
    %c4 = vm.const.i32 4 : i32
    vm.call.variadic @hal.command_buffer.execution_barrier(%ref_14, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
    vm.call @hal.command_buffer.end(%ref_14) : (!vm.ref<!hal.command_buffer>) -> ()
    vm.call @hal.ex.submit_and_wait(%ref, %ref_14) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
    %c16777224_23 = vm.const.i32 16777224 : i32
    %6 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224_23, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
    %c1_24 = vm.const.i32 1 : i32
    %7 = vm.call @hal.buffer.load(%ref_12, %6, %c1_24) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
    vm.cond_br %7, ^bb4(%2 : !vm.ref<!hal.buffer>), ^bb6(%2 : !vm.ref<!hal.buffer>)
  ^bb3:  // 2 preds: ^bb1, ^bb4
    %c2_25 = vm.const.i32 2 : i32
    vm.fail %c2_25, "unreachable location reached"
  ^bb4(%8: !vm.ref<!hal.buffer>):  // pred: ^bb2
    %c16777248_26 = vm.const.i32 16777248 : i32
    %9 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248_26) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
    %c50_27 = vm.const.i32 50 : i32
    %c15_28 = vm.const.i32 15 : i32
    %ref_29 = vm.call @hal.allocator.allocate(%ref_0, %c50_27, %c15_28, %9) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
    vm.call @hal.ex.defer_release(%ref_29) : (!vm.ref<!hal.buffer>) -> ()
    %c1_30 = vm.const.i32 1 : i32
    %c3_31 = vm.const.i32 3 : i32
    %ref_32 = vm.call @hal.command_buffer.create(%ref, %c1_30, %c3_31) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
    vm.call @hal.command_buffer.begin(%ref_32) : (!vm.ref<!hal.command_buffer>) -> ()
    %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
    %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
    %ref_33 = vm.call @hal.buffer.allocator(%8) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
    %c16777248_34 = vm.const.i32 16777248 : i32
    %10 = vm.call.variadic @hal.allocator.compute_size(%ref_33, [], %c16777248_34) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
    %zero_35 = vm.const.i32.zero : i32
    %zero_36 = vm.const.i32.zero : i32
    %c1_37 = vm.const.i32 1 : i32
    vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_32, %_executable_layout_1, %zero_35, [%zero_36, %c1_37], [%8, %ref_29], [%zero, %zero], [%10, %9]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
    %_device_match_id_0_38 = vm.global.load.i32 @_device_match_id_0 : i32
    vm.cond_br %_device_match_id_0_38, ^bb5, ^bb3
  ^bb5:  // pred: ^bb4
    %zero_39 = vm.const.i32.zero : i32
    vm.call @hal.command_buffer.dispatch(%ref_32, %_executable_while_ex_dispatch_1, %zero_39, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
    %c20_40 = vm.const.i32 20 : i32
    %c5_41 = vm.const.i32 5 : i32
    %c8_42 = vm.const.i32 8 : i32
    %c4_43 = vm.const.i32 4 : i32
    vm.call.variadic @hal.command_buffer.execution_barrier(%ref_32, %c20_40, %c5_41, [%c8_42], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
    vm.call @hal.command_buffer.end(%ref_32) : (!vm.ref<!hal.command_buffer>) -> ()
    vm.call @hal.ex.submit_and_wait(%ref, %ref_32) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
    vm.br ^bb1(%ref_29 : !vm.ref<!hal.buffer>)
  ^bb6(%11: !vm.ref<!hal.buffer>):  // pred: ^bb2
    %c16777248_44 = vm.const.i32 16777248 : i32
    %ref_45 = vm.call.variadic @hal.buffer_view.create(%11, [], %c16777248_44) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
    %c16777248_46 = vm.const.i32 16777248 : i32
    %ref_47 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248_46) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
    vm.call @check.expect_eq(%ref_45, %ref_47) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
    vm.return
  }
  vm.export @while as("while$raw")
  vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
    vm.call @while() : () -> ()
    vm.return
  }
  vm.export @while$sync as("while")
  vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
    %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
    vm.cond_fail %0, "semaphore wait failed"
    vm.call @while$sync() : () -> ()
    vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
    vm.return
  }
  vm.import @check.expect_true(%operand : i32) attributes {sym_visibility = "private"}
  vm.import @check.expect_false(%operand : i32) attributes {sym_visibility = "private"}
  vm.import @check.expect_all_true(%operand : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
  vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
  vm.import @check.expect_almost_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
  vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
  vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
  vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
  vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
  vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
  vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
  vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
  vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
  vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"}
  vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
  vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
  vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
  vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.buffer_view.trace(%operands : !vm.ref<!hal.buffer_view>...) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
  vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"}
  vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"}
  vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.semaphore.create(%device : !vm.ref<!hal.device>, %initial_value : i32) -> !vm.ref<!hal.semaphore> attributes {nosideeffects, sym_visibility = "private"}
  vm.import @hal.semaphore.query(%semaphore : !vm.ref<!hal.semaphore>) -> (i32, i32) attributes {sym_visibility = "private"}
  vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
  vm.import @hal.semaphore.fail(%semaphore : !vm.ref<!hal.semaphore>, %status : i32) attributes {sym_visibility = "private"}
  vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
  vm.func @__init() {
    %0 = vm.call @_device_match_id_0_initializer() : () -> i32
    vm.global.store.i32 %0, @_device_match_id_0 : i32
    %ref = vm.call @_descriptor_set_layout_0_initializer() : () -> !vm.ref<!hal.descriptor_set_layout>
    vm.global.store.ref %ref, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
    %ref_0 = vm.call @_executable_layout_0_initializer() : () -> !vm.ref<!hal.executable_layout>
    vm.global.store.ref %ref_0, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
    %ref_1 = vm.call @_descriptor_set_layout_1_initializer() : () -> !vm.ref<!hal.descriptor_set_layout>
    vm.global.store.ref %ref_1, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
    %ref_2 = vm.call @_executable_layout_1_initializer() : () -> !vm.ref<!hal.executable_layout>
    vm.global.store.ref %ref_2, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
    %ref_3 = vm.call @_executable_cache_initializer() : () -> !vm.ref<!hal.executable_cache>
    vm.global.store.ref %ref_3, @_executable_cache : !vm.ref<!hal.executable_cache>
    vm.return
  }
  vm.export @__init
}

// *** IR Dump After Inliner ***


module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.global.i32 @_device_match_id_0 mutable : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %c16777224 = vm.const.i32 16777224 : i32
      %c2 = vm.const.i32 2 : i32
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c3 = vm.const.i32 3 : i32
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %ref_2 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %ref_3 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
      %0 = iree.do_not_optimize(%ref_3) : !vm.ref<!hal.buffer>
      %1 = iree.do_not_optimize(%ref_2) : !vm.ref<!hal.buffer>
      vm.br ^bb1(%0 : !vm.ref<!hal.buffer>)
    ^bb1(%2: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_4 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %3) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_4) : (!vm.ref<!hal.buffer>) -> ()
      %ref_5 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_6 = vm.call @hal.buffer.allocator(%2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %4 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_7 = vm.call @hal.buffer.allocator(%1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %5 = vm.call.variadic @hal.allocator.compute_size(%ref_7, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_5, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%2, %1, %ref_4], [%zero, %zero, %zero], [%4, %5, %3]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      vm.call @hal.command_buffer.dispatch(%ref_5, %_executable_while_ex_dispatch_0, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_5, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_5) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %6 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %7 = vm.call @hal.buffer.load(%ref_4, %6, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %7, ^bb4(%2 : !vm.ref<!hal.buffer>), ^bb6(%2 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      vm.fail %c2, "unreachable location reached"
    ^bb4(%8: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %9 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_8 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %9) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_8) : (!vm.ref<!hal.buffer>) -> ()
      %ref_9 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_10 = vm.call @hal.buffer.allocator(%8) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %10 = vm.call.variadic @hal.allocator.compute_size(%ref_10, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_9, %_executable_layout_1, %zero, [%zero, %c1], [%8, %ref_8], [%zero, %zero], [%10, %9]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_11 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_11, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      vm.call @hal.command_buffer.dispatch(%ref_9, %_executable_while_ex_dispatch_1, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_9, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_9) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_8 : !vm.ref<!hal.buffer>)
    ^bb6(%11: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %ref_12 = vm.call.variadic @hal.buffer_view.create(%11, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %ref_13 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_12, %ref_13) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_br %0, ^bb2(%0 : i32), ^bb1
    ^bb1:  // pred: ^bb0
      vm.call @while() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    ^bb2(%1: i32):  // pred: ^bb0
      vm.fail %1, "semaphore wait failed"
    }
    vm.import @check.expect_true(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_false(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_all_true(%operand : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_almost_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.trace(%operands : !vm.ref<!hal.buffer_view>...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.create(%device : !vm.ref<!hal.device>, %initial_value : i32) -> !vm.ref<!hal.semaphore> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.query(%semaphore : !vm.ref<!hal.semaphore>) -> (i32, i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.fail(%semaphore : !vm.ref<!hal.semaphore>, %status : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
    vm.func @__init() {
      %c2 = vm.const.i32 2 : i32
      %c1 = vm.const.i32 1 : i32
      %c6 = vm.const.i32 6 : i32
      %zero = vm.const.i32.zero : i32
      %c1397773893 = vm.const.i32 1397773893 : i32
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.global.store.i32 %0, @_device_match_id_0 : i32
      %ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_1 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_0, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c1), (%c2, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_1, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref_2 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_3 = vm.call.variadic @hal.executable_layout.create(%ref_2, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_3, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_4 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_4, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_5, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref_6 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_7 = vm.call.variadic @hal.executable_layout.create(%ref_6, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_7, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_8 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_9 = vm.call @hal.executable_cache.create(%ref_8, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_10 = vm.switch.ref %1[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_11 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_0, %c7, %ref_10) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_11, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %2 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_12 = vm.switch.ref %2[%_while_ex_dispatch_1_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_13 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_1, %c7, %ref_12) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_13, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @_executable_cache : !vm.ref<!hal.executable_cache>
      vm.return
    }
    vm.export @__init
  }
}

// *** IR Dump After CSE ***


module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.global.i32 @_device_match_id_0 mutable : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %c16777224 = vm.const.i32 16777224 : i32
      %c2 = vm.const.i32 2 : i32
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c3 = vm.const.i32 3 : i32
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %ref_2 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %ref_3 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
      %0 = iree.do_not_optimize(%ref_3) : !vm.ref<!hal.buffer>
      %1 = iree.do_not_optimize(%ref_2) : !vm.ref<!hal.buffer>
      vm.br ^bb1(%0 : !vm.ref<!hal.buffer>)
    ^bb1(%2: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_4 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %3) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_4) : (!vm.ref<!hal.buffer>) -> ()
      %ref_5 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_6 = vm.call @hal.buffer.allocator(%2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %4 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_7 = vm.call @hal.buffer.allocator(%1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %5 = vm.call.variadic @hal.allocator.compute_size(%ref_7, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_5, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%2, %1, %ref_4], [%zero, %zero, %zero], [%4, %5, %3]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      vm.call @hal.command_buffer.dispatch(%ref_5, %_executable_while_ex_dispatch_0, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_5, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_5) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %6 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %7 = vm.call @hal.buffer.load(%ref_4, %6, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %7, ^bb4(%2 : !vm.ref<!hal.buffer>), ^bb6(%2 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      vm.fail %c2, "unreachable location reached"
    ^bb4(%8: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %9 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_8 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %9) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_8) : (!vm.ref<!hal.buffer>) -> ()
      %ref_9 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_10 = vm.call @hal.buffer.allocator(%8) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %10 = vm.call.variadic @hal.allocator.compute_size(%ref_10, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_9, %_executable_layout_1, %zero, [%zero, %c1], [%8, %ref_8], [%zero, %zero], [%10, %9]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_11 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_11, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      vm.call @hal.command_buffer.dispatch(%ref_9, %_executable_while_ex_dispatch_1, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_9, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_9) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_8 : !vm.ref<!hal.buffer>)
    ^bb6(%11: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %ref_12 = vm.call.variadic @hal.buffer_view.create(%11, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %ref_13 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_12, %ref_13) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_br %0, ^bb2(%0 : i32), ^bb1
    ^bb1:  // pred: ^bb0
      vm.call @while() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    ^bb2(%1: i32):  // pred: ^bb0
      vm.fail %1, "semaphore wait failed"
    }
    vm.import @check.expect_true(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_false(%operand : i32) attributes {sym_visibility = "private"}
    vm.import @check.expect_all_true(%operand : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @check.expect_almost_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_range(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.subspan(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.fill(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.read_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!iree.mutable_byte_buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.write_data(%target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %source_buffer : !vm.ref<!iree.byte_buffer>, %source_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.copy_data(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer.store(%value : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.subview(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.buffer(%buffer_view : !vm.ref<!hal.buffer_view>) -> !vm.ref<!hal.buffer> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.byte_length(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_offset(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.compute_range(%buffer_view : !vm.ref<!hal.buffer_view>, %indices : i32..., %lengths : i32...) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.rank(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dim(%buffer_view : !vm.ref<!hal.buffer_view>, %index : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.1(%buffer_view : !vm.ref<!hal.buffer_view>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.2(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.3(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.dims.4(%buffer_view : !vm.ref<!hal.buffer_view>) -> (i32, i32, i32, i32) attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.buffer_view.trace(%operands : !vm.ref<!hal.buffer_view>...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.fill_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32, %pattern : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.copy_buffer(%command_buffer : !vm.ref<!hal.command_buffer>, %source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %target_buffer : !vm.ref<!hal.buffer>, %target_offset : i32, %length : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_constants(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %offset : i32, %values : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.bind_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %descriptor_set : !vm.ref<!hal.descriptor_set>, %dynamic_offsets : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch.indirect(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroups_buffer : !vm.ref<!hal.buffer>, %workgroups_offset : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set.create(%device : !vm.ref<!hal.device>, %set_layout : !vm.ref<!hal.descriptor_set_layout>, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) -> !vm.ref<!hal.descriptor_set> attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.create(%device : !vm.ref<!hal.device>, %initial_value : i32) -> !vm.ref<!hal.semaphore> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.query(%semaphore : !vm.ref<!hal.semaphore>) -> (i32, i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.fail(%semaphore : !vm.ref<!hal.semaphore>, %status : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
    vm.func @__init() {
      %c2 = vm.const.i32 2 : i32
      %c1 = vm.const.i32 1 : i32
      %c6 = vm.const.i32 6 : i32
      %zero = vm.const.i32.zero : i32
      %c1397773893 = vm.const.i32 1397773893 : i32
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.global.store.i32 %0, @_device_match_id_0 : i32
      %ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_1 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_0, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c1), (%c2, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_1, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref_2 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_3 = vm.call.variadic @hal.executable_layout.create(%ref_2, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_3, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_4 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_4, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_5, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref_6 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_7 = vm.call.variadic @hal.executable_layout.create(%ref_6, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_7, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_8 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_9 = vm.call @hal.executable_cache.create(%ref_8, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_10 = vm.switch.ref %1[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_11 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_0, %c7, %ref_10) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_11, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %2 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_12 = vm.switch.ref %2[%_while_ex_dispatch_1_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_13 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_1, %c7, %ref_12) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_13, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @_executable_cache : !vm.ref<!hal.executable_cache>
      vm.return
    }
    vm.export @__init
  }
}

// *** IR Dump After SymbolDCE ***


module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.global.i32 @_device_match_id_0 mutable : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %c16777224 = vm.const.i32 16777224 : i32
      %c2 = vm.const.i32 2 : i32
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c3 = vm.const.i32 3 : i32
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %ref_2 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %ref_3 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
      %0 = iree.do_not_optimize(%ref_3) : !vm.ref<!hal.buffer>
      %1 = iree.do_not_optimize(%ref_2) : !vm.ref<!hal.buffer>
      vm.br ^bb1(%0 : !vm.ref<!hal.buffer>)
    ^bb1(%2: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_4 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %3) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_4) : (!vm.ref<!hal.buffer>) -> ()
      %ref_5 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_6 = vm.call @hal.buffer.allocator(%2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %4 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_7 = vm.call @hal.buffer.allocator(%1) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %5 = vm.call.variadic @hal.allocator.compute_size(%ref_7, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_5, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%2, %1, %ref_4], [%zero, %zero, %zero], [%4, %5, %3]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      vm.call @hal.command_buffer.dispatch(%ref_5, %_executable_while_ex_dispatch_0, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_5, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_5) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %6 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %7 = vm.call @hal.buffer.load(%ref_4, %6, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %7, ^bb4(%2 : !vm.ref<!hal.buffer>), ^bb6(%2 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      vm.fail %c2, "unreachable location reached"
    ^bb4(%8: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %9 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_8 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %9) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_8) : (!vm.ref<!hal.buffer>) -> ()
      %ref_9 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_10 = vm.call @hal.buffer.allocator(%8) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %10 = vm.call.variadic @hal.allocator.compute_size(%ref_10, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_9, %_executable_layout_1, %zero, [%zero, %c1], [%8, %ref_8], [%zero, %zero], [%10, %9]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_11 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_11, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      vm.call @hal.command_buffer.dispatch(%ref_9, %_executable_while_ex_dispatch_1, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_9, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_9) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_8 : !vm.ref<!hal.buffer>)
    ^bb6(%11: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %ref_12 = vm.call.variadic @hal.buffer_view.create(%11, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %ref_13 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_12, %ref_13) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_br %0, ^bb2(%0 : i32), ^bb1
    ^bb1:  // pred: ^bb0
      vm.call @while() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    ^bb2(%1: i32):  // pred: ^bb0
      vm.fail %1, "semaphore wait failed"
    }
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
    vm.func @__init() {
      %c2 = vm.const.i32 2 : i32
      %c1 = vm.const.i32 1 : i32
      %c6 = vm.const.i32 6 : i32
      %zero = vm.const.i32.zero : i32
      %c1397773893 = vm.const.i32 1397773893 : i32
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.global.store.i32 %0, @_device_match_id_0 : i32
      %ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_1 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_0, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c1), (%c2, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_1, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref_2 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_3 = vm.call.variadic @hal.executable_layout.create(%ref_2, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_3, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_4 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_4, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_5, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref_6 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_7 = vm.call.variadic @hal.executable_layout.create(%ref_6, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_7, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_8 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_9 = vm.call @hal.executable_cache.create(%ref_8, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_10 = vm.switch.ref %1[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_11 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_0, %c7, %ref_10) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_11, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %2 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_12 = vm.switch.ref %2[%_while_ex_dispatch_1_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_13 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_1, %c7, %ref_12) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_13, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @_executable_cache : !vm.ref<!hal.executable_cache>
      vm.return
    }
    vm.export @__init
  }
}

// *** IR Dump After mlir::iree_compiler::IREE::DropCompilerHintsPass ***


module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.global.i32 @_device_match_id_0 mutable : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %c16777224 = vm.const.i32 16777224 : i32
      %c2 = vm.const.i32 2 : i32
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c3 = vm.const.i32 3 : i32
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %ref_2 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %ref_3 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
      vm.br ^bb1(%ref_3 : !vm.ref<!hal.buffer>)
    ^bb1(%0: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %1 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_4 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_4) : (!vm.ref<!hal.buffer>) -> ()
      %ref_5 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_6 = vm.call @hal.buffer.allocator(%0) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %2 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_7 = vm.call @hal.buffer.allocator(%ref_2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_7, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_5, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%0, %ref_2, %ref_4], [%zero, %zero, %zero], [%2, %3, %1]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      vm.call @hal.command_buffer.dispatch(%ref_5, %_executable_while_ex_dispatch_0, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_5, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_5) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %4 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %5 = vm.call @hal.buffer.load(%ref_4, %4, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %5, ^bb4(%0 : !vm.ref<!hal.buffer>), ^bb6(%0 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      vm.fail %c2, "unreachable location reached"
    ^bb4(%6: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %7 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_8 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %7) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_8) : (!vm.ref<!hal.buffer>) -> ()
      %ref_9 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_10 = vm.call @hal.buffer.allocator(%6) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %8 = vm.call.variadic @hal.allocator.compute_size(%ref_10, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_9, %_executable_layout_1, %zero, [%zero, %c1], [%6, %ref_8], [%zero, %zero], [%8, %7]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_11 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_11, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      vm.call @hal.command_buffer.dispatch(%ref_9, %_executable_while_ex_dispatch_1, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_9, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_9) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_8 : !vm.ref<!hal.buffer>)
    ^bb6(%9: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %ref_12 = vm.call.variadic @hal.buffer_view.create(%9, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %ref_13 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_12, %ref_13) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_br %0, ^bb2(%0 : i32), ^bb1
    ^bb1:  // pred: ^bb0
      vm.call @while() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    ^bb2(%1: i32):  // pred: ^bb0
      vm.fail %1, "semaphore wait failed"
    }
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
    vm.func @__init() {
      %c2 = vm.const.i32 2 : i32
      %c1 = vm.const.i32 1 : i32
      %c6 = vm.const.i32 6 : i32
      %zero = vm.const.i32.zero : i32
      %c1397773893 = vm.const.i32 1397773893 : i32
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.global.store.i32 %0, @_device_match_id_0 : i32
      %ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_1 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_0, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c1), (%c2, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_1, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref_2 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_3 = vm.call.variadic @hal.executable_layout.create(%ref_2, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_3, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_4 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_4, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_5, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref_6 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_7 = vm.call.variadic @hal.executable_layout.create(%ref_6, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_7, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_8 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_9 = vm.call @hal.executable_cache.create(%ref_8, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_10 = vm.switch.ref %1[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_11 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_0, %c7, %ref_10) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_11, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %2 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_12 = vm.switch.ref %2[%_while_ex_dispatch_1_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_13 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_1, %c7, %ref_12) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_13, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @_executable_cache : !vm.ref<!hal.executable_cache>
      vm.return
    }
    vm.export @__init
  }
}



module {
  vm.module @module {
    vm.rodata @_utf8_vulkan_7197BF52A22CAFD7 dense<[118, 117, 108, 107, 97, 110, 42]> : vector<7xi8>
    vm.global.i32 @_device_match_id_0 mutable : i32
    vm.global.ref @_executable_while_ex_dispatch_0 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_executable_while_ex_dispatch_1 mutable : !vm.ref<!hal.executable>
    vm.global.ref @_descriptor_set_layout_0 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_0 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_descriptor_set_layout_1 mutable : !vm.ref<!hal.descriptor_set_layout>
    vm.global.ref @_executable_layout_1 mutable : !vm.ref<!hal.executable_layout>
    vm.global.ref @_executable_cache mutable : !vm.ref<!hal.executable_cache>
    vm.rodata @_utf8_default_7FD5254DFCA3A5D0 dense<[100, 101, 102, 97, 117, 108, 116]> : vector<7xi8>
    vm.rodata @_while_ex_dispatch_0_binary_spirv dense<"0x140000005350564500000A000C000000040008000A0000001C0400000400000004010000030223070000010016000000230000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000B0000007768696C655F65785F64697370617463685F3000100006000B0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F315F5F0000000005000800070000005F5F7265736F757263655F7661725F305F305F5F0000000005000800080000005F5F7265736F757263655F7661725F305F325F5F00000000050007000B0000007768696C655F65785F64697370617463685F3000470004000300000006000000040000004800050002000000000000002300000000000000470003000200000002000000470004000600000021000000010000004700040006000000220000000000000047000300020000000200000047000400070000002100000000000000470004000700000022000000000000004700030002000000020000004700040008000000210000000200000047000400080000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C0000003B00040001000000080000000C000000130002000A00000021000300090000000A0000002B000400040000000D000000000000002B000400040000000E000000080000002B000400040000000F000000FF0000002B00040004000000100000000400000020000400110000000C000000040000001400020016000000360005000A0000000B0000000000000009000000F80002000C000000410006001100000012000000070000000D0000000D0000003D000400040000001300000012000000410006001100000014000000060000000D0000000D0000003D000400040000001500000014000000B100050016000000170000001300000015000000A9000600040000001800000017000000050000000D0000008B00050004000000190000000D0000001000000084000500040000001A000000190000000E000000C4000500040000001B0000000F0000001A000000C8000400040000001C0000001B000000C7000500040000001D000000180000000F000000C4000500040000001E0000001D0000001A00000087000500040000001F0000000D00000010000000410006001100000020000000080000000D0000001F000000F0000700040000002100000020000000050000000E0000001C000000F1000700040000002200000020000000050000000E0000001E000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3000"> : vector<1108xi8>
    vm.rodata @_while_ex_dispatch_1_binary_spirv dense<"0x140000005350564500000A000C000000040008000A000000A402000004000000A6000000030223070000010016000000130000000000000011000200010000000A000B005350565F4B48525F73746F726167655F6275666665725F73746F726167655F636C617373000000000E00030000000000010000000F000800050000000A0000007768696C655F65785F64697370617463685F3100100006000A0000001100000001000000010000000100000005000800060000005F5F7265736F757263655F7661725F305F305F5F0000000005000800070000005F5F7265736F757263655F7661725F305F315F5F00000000050007000A0000007768696C655F65785F64697370617463685F310047000400030000000600000004000000480005000200000000000000230000000000000047000300020000000200000047000400060000002100000000000000470004000600000022000000000000004700030002000000020000004700040007000000210000000100000047000400070000002200000000000000150004000400000020000000000000002B0004000400000005000000010000001C0004000300000004000000050000001E000300020000000300000020000400010000000C000000020000003B00040001000000060000000C0000003B00040001000000070000000C00000013000200090000002100030008000000090000002B000400040000000C00000000000000200004000D0000000C0000000400000036000500090000000A0000000000000008000000F80002000B000000410006000D0000000E000000060000000C0000000C0000003D000400040000000F0000000E0000003D00040004000000100000000E0000008000050004000000110000000F00000010000000410006000D00000012000000070000000C0000000C0000003E0003001200000011000000FD000100380001000100000004000000130000007768696C655F65785F64697370617463685F3100"> : vector<732xi8>
    vm.rodata @while_const_0 dense<4> : tensor<i32>
    vm.rodata @while_const_1 dense<3> : tensor<i32>
    vm.rodata @while_const_2 dense<1> : tensor<i32>
    vm.func @while() {
      %c16777224 = vm.const.i32 16777224 : i32
      %c2 = vm.const.i32 2 : i32
      %c50 = vm.const.i32 50 : i32
      %c15 = vm.const.i32 15 : i32
      %c3 = vm.const.i32 3 : i32
      %c1 = vm.const.i32 1 : i32
      %zero = vm.const.i32.zero : i32
      %c20 = vm.const.i32 20 : i32
      %c5 = vm.const.i32 5 : i32
      %c8 = vm.const.i32 8 : i32
      %c4 = vm.const.i32 4 : i32
      %c16777248 = vm.const.i32 16777248 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_0 = vm.call @hal.device.allocator(%ref) : (!vm.ref<!hal.device>) -> !vm.ref<!hal.allocator>
      %while_const_0 = vm.const.ref.rodata @while_const_0 : !vm.ref<!iree.byte_buffer>
      %ref_1 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_0) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_1) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_1 = vm.const.ref.rodata @while_const_1 : !vm.ref<!iree.byte_buffer>
      %ref_2 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_1) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_2) : (!vm.ref<!hal.buffer>) -> ()
      %while_const_2 = vm.const.ref.rodata @while_const_2 : !vm.ref<!iree.byte_buffer>
      %ref_3 = vm.call.variadic @hal.allocator.allocate.const(%ref_0, %c50, %c15, [], %c16777248, %while_const_2) : (!vm.ref<!hal.allocator>, i32, i32, i32..., i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_3) : (!vm.ref<!hal.buffer>) -> ()
      vm.br ^bb1(%ref_3 : !vm.ref<!hal.buffer>)
    ^bb1(%0: !vm.ref<!hal.buffer>):  // 2 preds: ^bb0, ^bb5
      %1 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777224) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_4 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %1) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_4) : (!vm.ref<!hal.buffer>) -> ()
      %ref_5 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_0 = vm.global.load.ref @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_6 = vm.call @hal.buffer.allocator(%0) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %2 = vm.call.variadic @hal.allocator.compute_size(%ref_6, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_7 = vm.call @hal.buffer.allocator(%ref_2) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %3 = vm.call.variadic @hal.allocator.compute_size(%ref_7, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_5, %_executable_layout_0, %zero, [%zero, %c1, %c2], [%0, %ref_2, %ref_4], [%zero, %zero, %zero], [%2, %3, %1]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0, ^bb2, ^bb3
    ^bb2:  // pred: ^bb1
      vm.call @hal.command_buffer.dispatch(%ref_5, %_executable_while_ex_dispatch_0, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_5, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_5) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_5) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      %4 = vm.call.variadic @hal.allocator.compute_offset(%ref_0, [], %c16777224, []) : (!vm.ref<!hal.allocator>, i32..., i32, i32...) -> i32
      %5 = vm.call @hal.buffer.load(%ref_4, %4, %c1) : (!vm.ref<!hal.buffer>, i32, i32) -> i32
      vm.cond_br %5, ^bb4(%0 : !vm.ref<!hal.buffer>), ^bb6(%0 : !vm.ref<!hal.buffer>)
    ^bb3:  // 2 preds: ^bb1, ^bb4
      vm.fail %c2, "unreachable location reached"
    ^bb4(%6: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %7 = vm.call.variadic @hal.allocator.compute_size(%ref_0, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      %ref_8 = vm.call @hal.allocator.allocate(%ref_0, %c50, %c15, %7) : (!vm.ref<!hal.allocator>, i32, i32, i32) -> !vm.ref<!hal.buffer>
      vm.call @hal.ex.defer_release(%ref_8) : (!vm.ref<!hal.buffer>) -> ()
      %ref_9 = vm.call @hal.command_buffer.create(%ref, %c1, %c3) : (!vm.ref<!hal.device>, i32, i32) -> !vm.ref<!hal.command_buffer>
      vm.call @hal.command_buffer.begin(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      %_executable_while_ex_dispatch_1 = vm.global.load.ref @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_10 = vm.call @hal.buffer.allocator(%6) : (!vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator>
      %8 = vm.call.variadic @hal.allocator.compute_size(%ref_10, [], %c16777248) : (!vm.ref<!hal.allocator>, i32..., i32) -> i32
      vm.call.variadic @hal.command_buffer.push_descriptor_set(%ref_9, %_executable_layout_1, %zero, [%zero, %c1], [%6, %ref_8], [%zero, %zero], [%8, %7]) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable_layout>, i32, i32..., !vm.ref<!hal.buffer>..., i32..., i32...)
      %_device_match_id_0_11 = vm.global.load.i32 @_device_match_id_0 : i32
      vm.cond_br %_device_match_id_0_11, ^bb5, ^bb3
    ^bb5:  // pred: ^bb4
      vm.call @hal.command_buffer.dispatch(%ref_9, %_executable_while_ex_dispatch_1, %zero, %c1, %c1, %c1) : (!vm.ref<!hal.command_buffer>, !vm.ref<!hal.executable>, i32, i32, i32, i32) -> ()
      vm.call.variadic @hal.command_buffer.execution_barrier(%ref_9, %c20, %c5, [%c8], []) : (!vm.ref<!hal.command_buffer>, i32, i32, i32..., i32...)
      vm.call @hal.command_buffer.end(%ref_9) : (!vm.ref<!hal.command_buffer>) -> ()
      vm.call @hal.ex.submit_and_wait(%ref, %ref_9) : (!vm.ref<!hal.device>, !vm.ref<!hal.command_buffer>) -> ()
      vm.br ^bb1(%ref_8 : !vm.ref<!hal.buffer>)
    ^bb6(%9: !vm.ref<!hal.buffer>):  // pred: ^bb2
      %ref_12 = vm.call.variadic @hal.buffer_view.create(%9, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      %ref_13 = vm.call.variadic @hal.buffer_view.create(%ref_1, [], %c16777248) : (!vm.ref<!hal.buffer>, i32..., i32) -> !vm.ref<!hal.buffer_view>
      vm.call @check.expect_eq(%ref_12, %ref_13) : (!vm.ref<!hal.buffer_view>, !vm.ref<!hal.buffer_view>) -> ()
      vm.return
    }
    vm.export @while as("while$raw")
    vm.func @while$sync() attributes {iree.reflection = {f = "I1!R1!", fv = "1"}} {
      vm.call @while() : () -> ()
      vm.return
    }
    vm.export @while$sync as("while")
    vm.func @while$async(%arg0: !vm.ref<!hal.semaphore>, %arg1: i32, %arg2: !vm.ref<!hal.semaphore>, %arg3: i32) {
      %0 = vm.call @hal.semaphore.await(%arg0, %arg1) : (!vm.ref<!hal.semaphore>, i32) -> i32
      vm.cond_br %0, ^bb2(%0 : i32), ^bb1
    ^bb1:  // pred: ^bb0
      vm.call @while() : () -> ()
      vm.call @hal.semaphore.signal(%arg2, %arg3) : (!vm.ref<!hal.semaphore>, i32) -> ()
      vm.return
    ^bb2(%1: i32):  // pred: ^bb0
      vm.fail %1, "semaphore wait failed"
    }
    vm.import @check.expect_eq(%lhs : !vm.ref<!hal.buffer_view>, %rhs : !vm.ref<!hal.buffer_view>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.shared_device() -> !vm.ref<!hal.device> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.ex.defer_release(%operand : !vm.ref<?>) attributes {sym_visibility = "private"}
    vm.import @hal.ex.submit_and_wait(%device : !vm.ref<!hal.device>, %command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.allocator.compute_size(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.compute_offset(%allocator : !vm.ref<!hal.allocator>, %shape : i32..., %element_type : i32, %indices : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.allocator.allocate(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %allocation_size : i32) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.allocator.allocate.const(%allocator : !vm.ref<!hal.allocator>, %memory_types : i32, %buffer_usage : i32, %shape : i32..., %element_type : i32, %value : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.buffer> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.allocator(%buffer : !vm.ref<!hal.buffer>) -> !vm.ref<!hal.allocator> attributes {sym_visibility = "private"}
    vm.import @hal.buffer.load(%source_buffer : !vm.ref<!hal.buffer>, %source_offset : i32, %length : i32) -> i32 attributes {sym_visibility = "private"}
    vm.import @hal.buffer_view.create(%buffer : !vm.ref<!hal.buffer>, %shape : i32..., %element_type : i32) -> !vm.ref<!hal.buffer_view> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.command_buffer.create(%device : !vm.ref<!hal.device>, %modes : i32, %command_categories : i32) -> !vm.ref<!hal.command_buffer> attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.begin(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.end(%command_buffer : !vm.ref<!hal.command_buffer>) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.execution_barrier(%command_buffer : !vm.ref<!hal.command_buffer>, %source_stage_mask : i32, %target_stage_mask : i32, %memory_barriers : i32..., %buffer_barriers : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.push_descriptor_set(%command_buffer : !vm.ref<!hal.command_buffer>, %executable_layout : !vm.ref<!hal.executable_layout>, %set : i32, %bindings : i32..., %binding_buffers : !vm.ref<!hal.buffer>..., %binding_offsets : i32..., %binding_lengths : i32...) attributes {sym_visibility = "private"}
    vm.import @hal.command_buffer.dispatch(%command_buffer : !vm.ref<!hal.command_buffer>, %executable : !vm.ref<!hal.executable>, %entry_point : i32, %workgroup_x : i32, %workgroup_y : i32, %workgroup_z : i32) attributes {sym_visibility = "private"}
    vm.import @hal.descriptor_set_layout.create(%device : !vm.ref<!hal.device>, %usage_type : i32, %bindings : tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.allocator(%device : !vm.ref<!hal.device>) -> !vm.ref<!hal.allocator> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.device.match.id(%device : !vm.ref<!hal.device>, %pattern : !vm.ref<!iree.byte_buffer>) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.create(%device : !vm.ref<!hal.device>, %identifier : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.select_format(%executable_cache : !vm.ref<!hal.executable_cache>, %available_formats : i32...) -> i32 attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_cache.prepare(%executable_cache : !vm.ref<!hal.executable_cache>, %executable_layout : !vm.ref<!hal.executable_layout>, %caching_mode : i32, %executable_data : !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.executable_layout.create(%device : !vm.ref<!hal.device>, %set_layouts : !vm.ref<!hal.descriptor_set_layout>..., %push_constants : i32) -> !vm.ref<!hal.executable_layout> attributes {nosideeffects, sym_visibility = "private"}
    vm.import @hal.semaphore.signal(%semaphore : !vm.ref<!hal.semaphore>, %new_value : i32) attributes {sym_visibility = "private"}
    vm.import @hal.semaphore.await(%semaphore : !vm.ref<!hal.semaphore>, %min_value : i32) -> i32 attributes {sym_visibility = "private"}
    vm.func @__init() {
      %c2 = vm.const.i32 2 : i32
      %c1 = vm.const.i32 1 : i32
      %c6 = vm.const.i32 6 : i32
      %zero = vm.const.i32.zero : i32
      %c1397773893 = vm.const.i32 1397773893 : i32
      %null = vm.const.ref.zero : !vm.ref<!iree.byte_buffer>
      %c7 = vm.const.i32 7 : i32
      %ref = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_vulkan_7197BF52A22CAFD7 = vm.const.ref.rodata @_utf8_vulkan_7197BF52A22CAFD7 : !vm.ref<!iree.byte_buffer>
      %0 = vm.call @hal.device.match.id(%ref, %_utf8_vulkan_7197BF52A22CAFD7) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> i32
      vm.global.store.i32 %0, @_device_match_id_0 : i32
      %ref_0 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_1 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_0, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c1), (%c2, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_1, @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_0 = vm.global.load.ref @_descriptor_set_layout_0 : !vm.ref<!hal.descriptor_set_layout>
      %ref_2 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_3 = vm.call.variadic @hal.executable_layout.create(%ref_2, [%_descriptor_set_layout_0], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_3, @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %ref_4 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_5 = vm.call.variadic @hal.descriptor_set_layout.create(%ref_4, %c1, [(%zero, %c7, %c1), (%c1, %c7, %c6)]) : (!vm.ref<!hal.device>, i32, tuple<i32, i32, i32>...) -> !vm.ref<!hal.descriptor_set_layout>
      vm.global.store.ref %ref_5, @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %_descriptor_set_layout_1 = vm.global.load.ref @_descriptor_set_layout_1 : !vm.ref<!hal.descriptor_set_layout>
      %ref_6 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %ref_7 = vm.call.variadic @hal.executable_layout.create(%ref_6, [%_descriptor_set_layout_1], %zero) : (!vm.ref<!hal.device>, !vm.ref<!hal.descriptor_set_layout>..., i32) -> !vm.ref<!hal.executable_layout>
      vm.global.store.ref %ref_7, @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %ref_8 = vm.call @hal.ex.shared_device() : () -> !vm.ref<!hal.device>
      %_utf8_default_7FD5254DFCA3A5D0 = vm.const.ref.rodata @_utf8_default_7FD5254DFCA3A5D0 : !vm.ref<!iree.byte_buffer>
      %ref_9 = vm.call @hal.executable_cache.create(%ref_8, %_utf8_default_7FD5254DFCA3A5D0) : (!vm.ref<!hal.device>, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable_cache>
      %_executable_layout_0 = vm.global.load.ref @_executable_layout_0 : !vm.ref<!hal.executable_layout>
      %1 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_0_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_0_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_10 = vm.switch.ref %1[%_while_ex_dispatch_0_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_11 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_0, %c7, %ref_10) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_11, @_executable_while_ex_dispatch_0 : !vm.ref<!hal.executable>
      %_executable_layout_1 = vm.global.load.ref @_executable_layout_1 : !vm.ref<!hal.executable_layout>
      %2 = vm.call.variadic @hal.executable_cache.select_format(%ref_9, [%c1397773893]) : (!vm.ref<!hal.executable_cache>, i32...) -> i32
      %_while_ex_dispatch_1_binary_spirv = vm.const.ref.rodata @_while_ex_dispatch_1_binary_spirv : !vm.ref<!iree.byte_buffer>
      %ref_12 = vm.switch.ref %2[%_while_ex_dispatch_1_binary_spirv] else %null : !vm.ref<!iree.byte_buffer>
      %ref_13 = vm.call @hal.executable_cache.prepare(%ref_9, %_executable_layout_1, %c7, %ref_12) : (!vm.ref<!hal.executable_cache>, !vm.ref<!hal.executable_layout>, i32, !vm.ref<!iree.byte_buffer>) -> !vm.ref<!hal.executable>
      vm.global.store.ref %ref_13, @_executable_while_ex_dispatch_1 : !vm.ref<!hal.executable>
      vm.global.store.ref %ref_9, @_executable_cache : !vm.ref<!hal.executable_cache>
      vm.return
    }
    vm.export @__init
  }
}