ops::internal::Enter(root.WithOpName("enter_0"), a, "frame_0");
  Output exit_0 = ops::internal::Exit(root.WithOpName("exit_0"), enter_0);

  Output add = ops::Add(root.WithOpName("add"), exit_0, exit_0);

  Output enter_1 =
      ops::internal::Enter(root.WithOpName("enter_1"), add, "frame_0");
  Output exit_1 = ops::internal::Exit(root.WithOpName("exit_1"), enter_1);

  FixupSourceAndSinkEdges(root.graph());

  GraphCycles cycles;

// CHECK: %[[BROADCAST_1:.*]] = stablehlo.broadcast_in_dim %[[CONST_1]], dims = [3] : (tensor<8xf32>) -> tensor<1x1x2x8xf32>
// CHECK: %[[ADD:.*]] = stablehlo.add %[[MUL]], %[[BROADCAST_1]] : tensor<1x1x2x8xf32>
// CHECK: return %[[ADD]] : tensor<1x1x2x8xf32>

// -----

func.func @fuse_conv_batchnorm(%arg_0: tensor<1x3x4x3xf32>) -> (tensor<1x3x2x2xf32>) {

// conditions are also hoisted (i.e., `tf.Add` and `tf.Mul`).

// CHECK-LABEL: readvariableop_is_hoisted_if_readonly
// CHECK:       [[CST_0:%.*]] = "tf.Const"
// CHECK:       [[VAR:%.*]] = "tf.VarHandleOp"
// CHECK:       [[CST_1:%.*]] = "tf.Const"
// CHECK:       [[VAR_VAL:%.*]] = "tf.ReadVariableOp"
// CHECK:       [[RES_1:%.*]] = "tf.Add"([[VAR_VAL]], [[CST_0]])

  params_->cond_output = {less_than, 0};

  TF_Operation* one = ScalarConst(1, params_->body_graph, s_);
  TF_Operation* add =
      Add(params_->body_inputs[0], {one, 0}, params_->body_graph, s_);
  ASSERT_EQ(TF_OK, TF_GetCode(s_)) << TF_Message(s_);
  params_->body_outputs[0] = {add, 0};

  ExpectOK();

  // Create backprop graph
  TF_Output grad_output;

        "--add", "Microsoft.VisualStudio.Workload.VCTools", \
        "--add", "Microsoft.VisualStudio.Component.VC.Tools.x86.x64", \
        "--add", "Microsoft.VisualStudio.Component.Windows10SDK.19041" -Wait; \
    Start-Process -FilePath C:/TEMP/vs_buildtools.exe -ArgumentList "--installPath", "C:/VS", \
        "--quiet", "--wait", "--nocache", "--includeOptional", \

                                 Status run_status) {
    // Add FallbackEnabled pass that modifies the graph.
    auto optimization_pass =
        std::make_unique<NiceMock<ModifyMlirModulePass>>(run_status);
    ON_CALL(*optimization_pass, GetPassState(_, _, _, _))
        .WillByDefault(Return(pass_state));
    MlirOptimizationPassRegistry::Global().Add(10,
                                               std::move(optimization_pass));

    %mean: tensor<256xf32>, %variance: tensor<256xf32>)
    -> (tensor<4x256xf32>) {
  // CHECK-DAG: %[[EPS_BCAST:.+]] = mhlo.constant dense<1.001000e-05> : tensor<256xf32>
  // CHECK-DAG: %[[VARIANCE_EPS:.+]] = mhlo.add %[[VARIANCE]], %[[EPS_BCAST]] : tensor<256xf32>
  // CHECK-DAG: %[[VARIANCE_EPS_RSQRT:.+]] = mhlo.rsqrt %[[VARIANCE_EPS]] : tensor<256xf32>

  auto read_w = ops::ReadVariableOp(scope.WithOpName("ReadW"), arg6, DT_FLOAT);
  add_attrs(read_w.node());

  auto e = ops::Add(scope.WithOpName("E"), arg0, arg2);
  add_attrs(e.node());
  auto f = ops::Add(scope.WithOpName("F"), read_v, read_w);
  add_attrs(f.node());
  auto g = ops::Add(scope.WithOpName("G"), f, arg3);
  add_attrs(g.node());

  auto out0 = ops::_Retval(scope.WithOpName("b_identity_0_retval_RetVal"),

  RewritePatternSet pattern_list(&getContext());
  pattern_list.add<HandleOptionalFrom>(&getContext());
  pattern_list.add<HandleOptionalGet>(&getContext());
  pattern_list.add<HandleOptionalNone>(&getContext());
  pattern_list.add<HandleFunc>(&getContext());
  pattern_list.add<HandleCall>(&getContext());
  pattern_list.add<HandleIf>(&getContext());
  FrozenRewritePatternSet patterns(std::move(pattern_list));

  %1 = tfl.add %arg1, %arg1 {fused_activation_function = "NONE"} : tensor<*xf32>
  func.return %0, %1, %arg2 : tensor<*xi32>, tensor<*xf32>, tensor<i32>
}

// CHECK-LABEL: func private @WhileOp_cond(
// CHECK: tfl.greater
// CHECK-LABEL: func private @WhileOp_body(
// CHECK: tfl.sub
// CHECK: tfl.add

// -----