} : (tensor<1x4xf32>) -> (tensor<1x8xf32>)
  func.return %0 : tensor<1x8xf32>

  //       CHECK: %0 = "tfl.custom"(%arg0) <{
  //  CHECK-SAME:   custom_code = "custom_call.my_custom_op",
  //  CHECK-SAME:   custom_option = #tfl<const_bytes : "0x746869735F69735F615F746573745F737472696E67">
  //  CHECK-SAME: }> : (tensor<1x4xf32>) -> tensor<1x8xf32>
}

// CHECK-LABEL: mhlo_custom_call_test__dont_legalize_dict_backend_config

// CHECK-SAME: f = @composite_conv2d_fn_2}>
// CHECK-SAME: {_tfl_quant_trait = "fully_quantizable"
// CHECK: %[[BIASADD_0:.*]] = "tf.BiasAdd"(%[[PARTITIONEDCALL_0]], %[[CONST_0]])
// CHECK: %[[RELU6_0:.*]] = "tf.Relu6"(%[[BIASADD_0]])
// CHECK: %[[PARTITIONEDCALL_1:.*]] = "tf.PartitionedCall"(%arg0, %[[CONST_1]])
// CHECK-SAME: f = @composite_conv2d_fn_1

//CHECK: %0 = "vhlo.custom_call_v1"(%arg0, %arg1) <{
//CHECK-SAME: api_version = #vhlo<api_version_v1 API_VERSION_ORIGINAL>,
//CHECK-SAME: backend_config = #vhlo.string_v1<"">,
//CHECK-SAME: call_target_name = #vhlo.string_v1<"tf.ResizeBilinear">,
//CHECK-SAME: called_computations = #vhlo.array_v1<[]>,
//CHECK-SAME: has_side_effect = #vhlo.bool_v1<false>,
//CHECK-SAME: operand_layouts = #vhlo.array_v1<[]>,
//CHECK-SAME: output_operand_aliases = #vhlo.array_v1<[]>,

  #
  # CHECK:      func {{@[a-zA-Z_0-9]+}}(
  # CHECK-SAME:   %arg0: tensor<1xf32> {tf._user_specified_name = "d", tf_saved_model.index_path = [0, "a"]},
  # CHECK-SAME:   %arg1: tensor<2xf32> {tf._user_specified_name = "d", tf_saved_model.index_path = [0, "b"]},
  # CHECK-SAME:   %arg2: tensor<3xf32> {tf._user_specified_name = "d", tf_saved_model.index_path = [0, "c"]},

// CHECK:           "mhlo.dot_general"({{.*}}, {{.*}}) <{
// CHECK-SAME:        lhs_batching_dimensions = [0]
// CHECK-SAME:        rhs_batching_dimensions = [0]
// CHECK-SAME:        lhs_contracting_dimensions = [2]
// CHECK-SAME:        rhs_contracting_dimensions = [1]

  // CHECK: %[[CONV2D_BACKPROP:[0-9]*]] = "tf.Conv2DBackpropFilter"
  // CHECK-SAME: (%[[IN_TRANSPOSE]], %[[FILTER_PERM]], %[[OUT_BP_TRANSPOSE]])
  // CHECK-SAME: data_format = "NCHW"
  // CHECK-SAME: dilations = [1, 4, 2, 3]
  // CHECK-SAME: explicit_paddings = [1, 2, 7, 8, 3, 4, 5, 6]
  // CHECK-SAME: padding = "EXPLICIT"
  // CHECK-SAME: strides = [5, 8, 6, 7]
  // CHECK-SAME: (tensor<1x3x32x32xf32>, tensor<4xi32>, tensor<1x8x32x32xf32>)

    // CHECK-SAME: device = "/job:localhost/replica:0/task:0/device:TPU:0"
    %1:2 = tf_executor.Enter %arg0 frame "frame" : tensor<i64>
    tf_executor.fetch %1#0 : tensor<i64>
  }
  func.return %0 : tensor<i64>
}

// CHECK-LABEL: func @testExitOp
// CHECK-SAME: ({{%.+}}: tensor<i64> {tf.device = "/job:localhost/replica:0/task:0/device:TPU:0"})

// CHECK-SAME: value = dense<0> : tensor<3x2xi32>
// CHECK-DAG: %[[LHS_DILATION:.*]] = "tf.Const"()
// CHECK-SAME: value = dense<[4, 1, 1]> : tensor<3xi32>
// CHECK-DAG: %[[LHS_DILATION_COPY:.*]] = "tf.Const"()
// CHECK-SAME: value = dense<[4, 1, 1]> : tensor<3xi32>
// CHECK-DAG: %[[RHS_DILATION:.*]] = "tf.Const"()
// CHECK-SAME: value = dense<1> : tensor<3xi32>
// CHECK-DAG: %[[RHS_DILATION_COPY:.*]] = "tf.Const"()

# it is being invoked.
# CHECK: module
# CHECK-SAME: tf.versions
# CHECK-SAME: bad_consumers
# CHECK-SAME: min_consumer
# CHECK-SAME: producer

# CHECK: tf_saved_model.session_initializer
# CHECK-SAME: initializers = [@[[restore:.*]]]

# CHECK: "tf_saved_model.asset"()
# CHECK-SAME: <{filename = [[filename:.*]], sym_name = "[[sym_name:.*]]"}> : () -> ()

# UNKNOWN-SAME:  outputs = "Add"
# UNKNOWN:           %[[add:.*]], %[[add_control:.*]] = tf_executor.island wraps "tf.Add"(%[[ARG_0]], %[[ARG_1]])
# UNKNOWN:           fetch %[[add]]

# DYNAMIC-LABEL: func @main
# DYNAMIC-SAME:  (%[[ARG_0:[a-z0-9]+]]: tensor<?x1x?xi32>, %[[ARG_1:[a-z0-9]+]]: tensor<1x?x1xi32>) -> tensor<*xi32>
# DYNAMIC-SAME:  control_outputs = ""
# DYNAMIC-SAME:  inputs = "input0,input1"