return %0 : tensor<1x3xf32>
}

// CHECK-LABEL: @_func
func.func private @_func(%arg0: tensor<1x3xf32>) -> (tensor<1x3xf32>) {
  return %arg0 : tensor<1x3xf32>
}
}


// -----
// SPMD: one input and no return
//
// CHECK-LABEL: func.func @serving_default(%arg0: tensor<1x3xf32>) {
// CHECK-NEXT:  "tf.IfrtCall"(%arg0)
// CHECK-SAME:       {program_id = [[PROGRAM_ID:.*]] : i64, variable_arg_indices = []}

		{name: "CMPU", argLength: 2, reg: gp2flags, asm: "CMPU", typ: "Flags"},   // arg0 compare to arg1
		{name: "CMPWU", argLength: 2, reg: gp2flags, asm: "CMPWU", typ: "Flags"}, // arg0 compare to arg1

		{name: "CMPconst", argLength: 1, reg: gp1flags, asm: "CMP", typ: "Flags", aux: "Int32"},     // arg0 compare to auxint

// CHECK-NOT: tf._noinline
func.func @inner_stateful_pcall_func(%arg0: tensor<i32>) -> tensor<i32> {
  %0 = "tf.StatefulPartitionedCall"(%arg0) {_xla_compile_device_type = "CPU", config = "", config_proto = "", device = "/device:CPU:0", executor_type = "", f = @func} : (tensor<i32>) -> (tensor<i32>)
  func.return %0 : tensor<i32>
}

func.func @func(%arg0: tensor<i32>) -> tensor<i32> {
  func.return %arg0 : tensor<i32>

// CHECK:           return %[[VAL_3]] : tensor<2xf32>
// CHECK:         }
func.func @floordiv_f32(%arg0: tensor<2xf32>) -> tensor<2xf32> {
  %0 = mhlo.divide %arg0, %arg0 : tensor<2xf32>
  %1 = mhlo.divide %arg0, %arg0 : tensor<2xf32>
  %2 = "mhlo.floor"(%1) : (tensor<2xf32>) -> tensor<2xf32>
  func.return %2 : tensor<2xf32>
}

// CHECK-LABEL:   func @floordiv_f16_broadcast(

// CHECK-LABEL: func private @batched_function
func.func private @batched_function(%arg0: tensor<1x3xf32>) -> tensor<1x3xf32> {
  %2 = "tf.Identity"(%arg0) : (tensor<1x3xf32>) -> tensor<1x3xf32>
  func.return %2 : tensor<1x3xf32>
}

// CHECK-LABEL: func @main
func.func @main(%arg0: tensor<1x3xf32>) -> tensor<*xf32> {
  // CHECK:  "tf.BatchFunction"
  // CHECK-SAME: allowed_batch_sizes = [6]

// CHECK-LABEL: nchw_conv
// CHECK-SAME: %[[ARG:.+]]: tensor<1x8x4x4xf32>
func.func @nchw_conv(%arg0: tensor<1x8x4x4xf32>) -> tensor<1x8x4x4xf32> {
  %0 = stablehlo.constant() {value = dense<7.000000e+00> : tensor<8x8x3x3xf32>} : () -> tensor<8x8x3x3xf32>

// CHECK-LABEL: func @test_supported_lowering_of_tf_if_region1
// CHECK-SAME: (%[[ARG0:.*]]: tensor<i1>, %[[ARG1:.*]]: tensor<4xf32>)
func.func @test_supported_lowering_of_tf_if_region1(%arg0: tensor<i1>, %arg1: tensor<4xf32>) -> (tensor<*xf32>, tensor<4xf32>) {
  %res:2 = "tf.IfRegion"(%arg0) ({
    %call = func.call @test_if_then1(%arg1) : (tensor<4xf32>) -> tensor<4xf32>

// CHECK-LABEL: @main
func.func @main(%arg0: tensor<1xf32>) -> tensor<1xf32> {
  func.return %arg0: tensor<1xf32>
}
// CHECK-NEXT: return

// CHECK-LABEL: @unneeded_control_node_gets_deleted
func.func @unneeded_control_node_gets_deleted(%arg0: tensor<1xf32>)->tensor<1xf32> {
  %tmp0, %ctl0  = tfl.control_node controls "tfl.neg"(%arg0): (tensor<1xf32>) -> tensor<1xf32>
  func.return %tmp0: tensor<1xf32>
}

func libfuzzerTraceCmp1(arg0, arg1 uint8, fakePC uint)  {}
func libfuzzerTraceCmp2(arg0, arg1 uint16, fakePC uint) {}
func libfuzzerTraceCmp4(arg0, arg1 uint32, fakePC uint) {}
func libfuzzerTraceCmp8(arg0, arg1 uint64, fakePC uint) {}

func libfuzzerTraceConstCmp1(arg0, arg1 uint8, fakePC uint)  {}
func libfuzzerTraceConstCmp2(arg0, arg1 uint16, fakePC uint) {}

// CHECK-SAME: ([[ARG0:%.*]]: tensor<!tf_type.resource<tensor<10x3xf32>>>)
func.func @simple_packed(%arg0: tensor<!tf_type.resource<tensor<10x3xf32>>>) -> tensor<!tf_type.resource<tensor<10x3xf32>>> {
  // CHECK: "tf.TPUReplicateMetadata"()
  // CHECK: [[RI_0:%.*]] = "tf.TPUReplicatedInput"([[ARG0]])
  // CHECK-SAME: is_packed = true
  // CHECK: [[RI_1:%.*]] = "tf.TPUReplicatedInput"([[ARG0]])
  // CHECK-SAME: is_packed = true