%3 = "tf.While"(%1) {body = @while_body, cond = @while_cond}
        tf_device.return %3 : tensor<f32>
      }) : () -> tensor<f32>
      "tf.AssignVariableOp"(%0, %2)
      return
    }
    func @while_body(%arg0: tensor<f32>) {
      %0 = "tf.Const"() ...
      return %0 : tensor<f32>
    }
    func @while_cond(%arg0: tensor<f32>) {
      return %arg0
    }
    ```
  }];

  // CHECK: device = "xpu"
  %1 = "tf.MatMul"(%arg0, %0) {T = f32, _output_shapes = ["tfshape$dim { size: 3 } dim { size: 3 }"], device = "", transpose_a = false, transpose_b = false} : (tensor<3x1xf32>, tensor<1x3xf32>) -> tensor<3x3xf32>
  // CHECK: device = "cpu"
  %2 = "tf.Relu"(%1) {T = f32, _output_shapes = ["tfshape$dim { size: 3 } dim { size: 3 }"], device = "cpu"} : (tensor<3x3xf32>) -> tensor<3x3xf32>

  // CHECK-NEXT: [[r2:%.*]] = tfrt_fallback_async.executeop {{.*}} "tf.Tanh"([[r1]]) {T = f32}
  %4 = "tf.Tanh"(%3) {T = f32, _output_shapes = ["tfshape$dim { size: 3 } dim { size: 3 }"], device = "/device:CPU:0"} : (tensor<3x3xf32>) -> tensor<3x3xf32>
  // CHECK-NOT: tf.Identity
  %5 = "tf.Identity"(%4) {T = f32, _output_shapes = ["tfshape$dim { size: 3 } dim { size: 3 }"], device = "/device:CPU:0"} : (tensor<3x3xf32>) -> tensor<3x3xf32>

  %0 = "tf.FakeQuantWithMinMaxVars"(%arg0, %arg1, %arg2) : (tensor<8x8x8x8xf32>, tensor<f32>, tensor<f32>) -> tensor<8x8x8x8xf32>
  func.return %0 : tensor<8x8x8x8xf32>
}

// -----

// Test invalid tf.FakeQuantWithMinMaxVars
func.func @testInvalidFakeQuantWithMinMaxVarsWrongAttr(tensor<8x8x8x8xf32>, tensor<f32>, tensor<f32>) -> tensor<8x8x8x8xf32> {

}

// CHECK-LABEL: @splat_scalar_no_change
func.func @splat_scalar_no_change() -> (tensor<f32>, tensor<i32>) {
  // CHECK-NOT: mhlo.broadcast_in_dim
  %cst0 = mhlo.constant dense<0.000000e+00> : tensor<f32>
  %cst1 = mhlo.constant dense<0> : tensor<i32>
  func.return %cst0, %cst1 : tensor<f32>, tensor<i32>

        %3 = "tf.opA"(%arg2) : (tensor<i1>) -> tensor<f32>
        %4 = "tf.opB"(%arg2) : (tensor<i1>) -> tensor<i32>
        tf_device.return %3, %4 : tensor<f32>, tensor<i32>
      }
      tf_executor.yield %2#0, %2#1, %2#2, %2#3 : tensor<f32>, tensor<f32>, tensor<i32>, tensor<i32>
    }
    tf_executor.fetch %1#0, %1#1, %1#2, %1#3 : tensor<f32>, tensor<f32>, tensor<i32>, tensor<i32>
  }
  func.return
}

    %b = builtin.unrealized_conversion_cast to tensor<f32>
    %0, %1, %2 = "tf.WhileRegion"(%a, %b, %arg0) ({
      ^bb0(%ca: tensor<f32>, %cb: tensor<f32>, %cc: !res):
      %cond = builtin.unrealized_conversion_cast to tensor<i1>
      "tf.Yield"(%cond, %ca, %cb, %cc) : (tensor<i1>, tensor<f32>, tensor<f32>, !res) -> ()
    }, {
      ^bb0(%ba: tensor<f32>, %bb: tensor<f32>, %bc: !res):

// 8-bit integer and expressed type is f32.
bool IsI8F32UniformQuantizedType(Type type);

// Returns true iff `type` is a uniform quantized per-axis (per-channel) type
// whose storage type is 8-bit integer and expressed type is f32.
bool IsI8F32UniformQuantizedPerAxisType(Type type);

// Returns true iff `type` is a uniform quantized type whose storage type is
// 32-bit integer and expressed type is f32.

  // CHECK-NEXT:  %[[v1:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f32>
  // CHECK-NEXT:  %[[v2:.*]] = mhlo.constant dense<2.000000e+00> : tensor<f32>
  // CHECK-NEXT:  %[[v3:.*]] = mhlo.constant dense<-1.280000e+02> : tensor<f32>
  // CHECK-NEXT:  %[[v4:.*]] = "mhlo.broadcast_in_dim"(%3) <{broadcast_dimensions = dense<> : tensor<0xi64>}> : (tensor<f32>) -> tensor<1x300x300x40xf32>

  %0 = corert.const_dense_tensor dense<0.0> : tensor<f32>
  // CHECK: tfrt_fallback_async.corert_tensorhandle_to_fallback_tensor
  %1 = tfrt_fallback_async.corert_tensorhandle_to_fallback_tensor %0 {_tfrt_cost = 1 : i64, device = "cpu"} : (!corert.tensorhandle) -> (!tfrt_fallback.tf_tensor)
  // CHECK: tfrt_fallback_async.executeop key(0) cost(100) device("cpu") "tf.Relu"(%{{.*}}) {T = f32} : 1