// CHECK-LABEL @export_three
module @export_three attributes {tf_saved_model.semantics} {
  // CHECK: func @a
  // CHECK-SAME: i32
  // CHECK-SAME: f32
  // CHECK-SAME: tensor<i32>
  // CHECK: [[tuple:%.*]]:2 = call @f
  // CHECK: return [[tuple]]#0, [[tuple]]#1
  // CHECK-SAME: tensor<i8>
  // CHECK-SAME: tensor<i16>
  // CHECK: func @b
  // CHECK: func @c
  // CHECK: func private @f
  // CHECK-NOT: exported_names

  // CHECK: "tf.XlaHostCompute"(%arg0, %arg1)
  // CHECK-SAME: ancestors = []
  // CHECK-SAME: cost_estimate_ns = 1000000 : i64
  // CHECK-SAME: key = ""
  // CHECK-SAME: recv_key = "host_compute_channel_recv"
  // CHECK-SAME: send_key = "host_compute_channel_send"
  // CHECK-SAME: shape_inference_graph = @host_func
  // CHECK-SAME: shapes = [#tf_type.shape<*>, #tf_type.shape<3x?>]
  // CHECK-SAME: tpu_core = 0 : i64
  // CHECK: func @host_func

  // CHECK-SAME: ([[serialized]]
  // CHECK-NOT: device
  // CHECK-SAME: Tdense = [f32, f32]
  // CHECK-SAME: dense_shapes = [#corert.shape<>, #corert.shape<>]
  // CHECK-SAME: num_sparse = 2 : i64
  // CHECK-SAME: ragged_split_types = []
  // CHECK-SAME: ragged_value_types = []
  // CHECK-SAME: sparse_types = [!corert.string, i64]

  func.return %0 : tensor<i32>
}

// CHECK-LABEL: func private @merge_stateless_merged_if_0_0_then
// CHECK-SAME: ([[x:%.*]]: tensor<i32>, [[y:%.*]]: tensor<i32>)
// CHECK: [[r0:%.*]] = "tf.AddV2"([[x]], [[y]])
// CHECK: return [[r0]], [[r0]], [[r0]]

// CHECK-LABEL: func private @merge_stateless_merged_if_0_0_else
// CHECK-SAME: ([[x:%.*]]: tensor<i32>, [[y:%.*]]: tensor<i32>)

  // CHECK: %[[CONST:.*]] = mhlo.constant()
  // CHECK-SAME{LITERAL} value = dense<127> : tensor<2x3x3x2xi8>
  // CHECK-SAME: tensor<2x3x3x2x!quant.uniform<i8:f32, 1.000000e+00:3>>
  // CHECK: mhlo.convolution(%arg0, %[[CONST]])
  // CHECK-SAME{LITERAL}: dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]
  // CHECK-SAME{LITERAL}: window = {stride = [1, 2], pad = [[0, 0], [0, 0]], lhs_dilate = [1, 1], rhs_dilate = [2, 2]}

  // CHECK: "tf_saved_model.session_initializer"()
  // CHECK-SAME: initializers = [@init_func_restore_op]

  func.func @init_func_restore_op() -> () attributes {
      tf_saved_model.initializer_type = "restore_op",
      tf_saved_model.exported_names = ["tf_saved_model.session_initializer_restore_op"]} {
    return
  }
  // CHECK: @init_func_restore_op
  // CHECK-SAME: tf_saved_model.initializer_type = "restore_op"

  %0 = "tf_device.cluster"() ({
    // CHECK: "tf.Const"() <{value = dense<1> : tensor<i32>}>
    // CHECK-NOT: _xla_outside_compilation
    // CHECK: "tf.Const"
    // CHECK-SAME: _xla_outside_compilation
    // CHECK-SAME: tf_type.string
    // CHECK: "tf.Identity"
    // CHECK-NOT: _xla_outside_compilation
    %1 = "tf.Const"() {value = dense<1> : tensor<i32>} : () -> tensor<i32>

  // CHECK:      "tf.EnqueueTPUEmbeddingArbitraryTensorBatch"
  // CHECK-SAME:   device_ordinal = 0
  // CHECK:      "tf.EnqueueTPUEmbeddingArbitraryTensorBatch"
  // CHECK-SAME:   device_ordinal = 1
  // CHECK:      "tf.EnqueueTPUEmbeddingArbitraryTensorBatch"
  // CHECK-SAME:   device_ordinal = 0
  // CHECK:      "tf.EnqueueTPUEmbeddingArbitraryTensorBatch"
  // CHECK-SAME:   device_ordinal = 1
}

// -----

// CHECK: {tf_saved_model.index_path = ["input2"]}
// CHECK-SAME: {tf_saved_model.index_path = ["input1"]}
// CHECK-SAME: {tf_saved_model.index_path = ["start_logits"]}
// CHECK-SAME: {tf_saved_model.index_path = ["end_logits"]}
// CHECK-SAME: tf.entry_function = {inputs = "serving_default_input2:0,serving_default_input1:0", outputs = "StatefulPartitionedCall:1,StatefulPartitionedCall:0"}
// CHECK-SAME: tf_saved_model.exported_names = ["serving_default"]

    // CHECK-SAME:       device_ordinal = 0
    // CHECK-SAME:       device_type = "TPU"
    // CHECK-SAME:       key = "host_compute_channel_0_retvals"
    // CHECK:          "tf_device.cluster"
    // CHECK:            %[[A_OUTPUT:[0-9]*]] = "tf.A"
    // CHECK:            %[[HOST_OUTPUT:[0-9]*]] = "tf._XlaHostComputeMlir"()
    // CHECK-SAME:       recv_key = "host_compute_channel_0_retvals"