func.return
}

// CHECK-LABEL: func @device_func_5
// CHECK-SAME: [[ARG0:%.*]]: tensor<f32> {tf.aliasing_output = 1 : i64},
// CHECK-SAME: [[ARG1:%.*]]: tensor<f32> {tf.aliasing_output = 2 : i64},
// CHECK-NOT: [[ARG2:%.*]]:  tensor<f32> {tf.aliasing_output = 0 : i64}
func.func @device_func_5(%arg0: tensor<f32>, %arg1: tensor<f32>, %arg2: tensor<f32>) -> (tensor<f32>, tensor<f32>, tensor<f32>) {

func.func @no_operand_replicate() {
  tf_device.replicate {n = 2 : i32} {
    %0 = "tf.Const"() { value = dense<0> : tensor<i64> } : () -> tensor<i64>
    %1 = "tf.Const"() { value = dense<1> : tensor<i64> } : () -> tensor<i64>
    tf_device.return %0, %1 : tensor<i64>, tensor<i64>
  }
  func.return
  // CHECK:      tf_device.replicate
  // CHECK-SAME: n = 2
  // CHECK:   tf_device.return
}

// -----

// CHECK-SAME: Tdense = [f32, f32], dense_shapes = [#corert.shape<>, #corert.shape<>]
// CHECK-SAME: num_sparse = 2 : i64, ragged_split_types = [], ragged_value_types = []
// CHECK-SAME: sparse_types = [!corert.string, i64]}
// CHECK-SAME: num_args(7)

// CHECK: tfrt_fallback_async.createop(%0) key(1) device("/device:CPU:0") "tf.ReadVariableOp"() {dtype = f32} num_args(1)

      window = {
        stride = [1, 1], pad = [[0, 0], [1, 1]],
        lhs_dilate = [1, 1],
        rhs_dilate = [1, 1]
      }
      {
        batch_group_count = 1 : i64,
        feature_group_count = 1 : i64
      } : (tensor<1x3x2x3xf32>, tensor<2x3x3x2xf32>)
      -> tensor<1x2x2x2xf32>
    %3 = "quantfork.stats"(%arg2) {layerStats = dense<[7.05456924, 7.11401462]> : tensor<2xf32>} : (tensor<2xf32>) -> tensor<2xf32>

      %1 = "quantfork.qcast"(%0) {volatile} : (tensor<1x3xf32>) -> tensor<1x3x!quant.uniform<i8:f32, 0.13170163023705575:-1>>

    %0 = stablehlo.convolution(%arg0, %arg1) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3xf32>, tensor<2x3x3x2xf32>) -> tensor<1x3x4x2xf32>
    return %0 : tensor<1x3x4x2xf32>
  }
}

// CHECK-LABEL: quantize_conv_fn
// CHECK-SAME: %[[ARG0:.+]]: tensor<1x3x4x3xf32>

      %outputs, %control = tf_executor.island {
        "tf.TPUReplicateMetadata"() <{allow_soft_placement = true, computation_shape = [], device_assignment = [], host_compute_core = [], num_cores_per_replica = 1 : i64, num_replicas = 2 : i64, padding_map = [], step_marker_location = "STEP_MARK_AT_ENTRY", topology = "", tpu_compile_options_proto = "", use_spmd_for_xla_partitioning = false, use_tpu = true}> {_tpu_replicate = "cluster", device = ""} : () -> ()

    %3 = "tf.XlaCallModule"(%2) <{Sout = [#tf_type.shape<6x5>], dim_args_spec = [], disabled_checks = [], has_token_input_output = false, module = "", platforms = ["CPU"], version = 9 : i64}> {_entry_function = @main_1, _stablehlo_module_attrs = {}, device = ""} : (tensor<3x10xf32>)...

  %conv = stablehlo.convolution(%quant_arg, %quant_cst) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3x!quant.uniform<i8:f32, 0.0039207626791561354:-128>>, tensor<2x3x3x2x!quant.uniform<i8<-127:127>:f32, 0.015>>) -> tensor<1x3x4x2x!quant.uniform<i32:f32, 5.8949912267181218E-5>>

          "Failed at tf.StridedSlice op\nsee current operation: %2 = "
          "\"tf.StridedSlice\"(%arg0, %1, %1, %0) <{begin_mask = 11 : "
          "i64, ellipsis_mask = 0 : i64, end_mask = 11 : i64, new_axis_mask = "
          "4 : i64, shrink_axis_mask = 0 : i64}> {device = \"\"} : "
          "(tensor<*xf32>, tensor<4xi32>, tensor<4xi32>, tensor<4xi32>) "
          "-> tensor<*xf32>\nError code: ERROR_NEEDS_FLEX_OPS",