// CHECK-SAME: channel_handle = #mhlo.channel_handle
  // CHECK-SAME: handle =
  // CHECK-SAME: type = 3
  // CHECK-SAME: is_host_transfer = true
  // CHECK-SAME: mhlo.frontend_attributes = {_xla_host_transfer_handler_name = "tf_rendezvous", _xla_host_transfer_rendezvous = "recv_key_htod_0"}
  // CHECK-SAME: mhlo.sharding = "\08\01\1A\01\01\22\01\00"
  // CHECK-SAME: (!mhlo.token) -> (tensor<i32>, !mhlo.token)

// CHECK:      %[[RES:.*]] = "mhlo.dot_general"(%arg0, %arg1) <{
// CHECK-SAME:   dot_dimension_numbers = #mhlo.dot<
// CHECK-SAME:     lhs_batching_dimensions = [0],
// CHECK-SAME:     rhs_batching_dimensions = [0],
// CHECK-SAME:     lhs_contracting_dimensions = [3],
// CHECK-SAME:     rhs_contracting_dimensions = [1]
// CHECK-SAME: >}> : (tensor<3x72x1x2048xf32>, tensor<3x2048x512xf32>) -> tensor<3x72x1x512xf32>

// CHECK-SAME: batch_group_count = 1 : i64
// CHECK-SAME: dimension_numbers = "\10\03\1A\02\01\02 \02(\032\02\00\01@\03J\02\01\02"
// CHECK-SAME: explicit_padding = []
// CHECK-SAME: feature_group_count = 1 : i64
// CHECK-SAME: lhs_dilation = [1, 1]
// CHECK-SAME: padding = "VALID"
// CHECK-SAME: rhs_dilation = [2, 2]
// CHECK-SAME: rhs_quantization_axis = -1 : i64

// 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-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"()

// CHECK-SAME: f = @composite_conv2d_with_bias_and_relu6_fn_1}>
// CHECK-SAME: {_tfl_quant_trait = "fully_quantizable"
// CHECK: %[[PARTITIONEDCALL_1:.*]] = "tf.PartitionedCall"(%arg0, %arg1, %[[CONST_0]])
// CHECK-SAME: f = @composite_conv2d_with_bias_and_relu_fn_1}
// CHECK: %[[PARTITIONEDCALL_2:.*]] = "tf.PartitionedCall"(%arg0, %arg1, %[[CONST_0]])
// CHECK-SAME: f = @composite_conv2d_with_bias_fn_1}

// CHECK-LABEL: func @interfering_accesses
// CHECK-SAME: %[[ARG_0:.*]]: tensor<*x!tf_type.resource<tensor<32xf32>>>
// CHECK-SAME: %[[ARG_1:.*]]: tensor<*x!tf_type.resource<tensor<64xf32>>>
// CHECK-SAME: %[[ARG_2:.*]]: tensor<32xf32>
// CHECK-SAME: %[[ARG_4:.*]]: tensor<*x!tf_type.resource<tensor<8xf32>>>
// CHECK-SAME: %[[ARG_5:.*]]: tensor<*x!tf_type.resource<tensor<2xf32>>>
// CHECK-SAME: %[[ARG_6:.*]]: tensor<2xf32>

        : () -> !tf_res
      // CHECK-NEXT: "tf.Identity"
      // CHECK-SAME: {device = "/TPU:0"}
      %id0 = "tf.Identity"(%arg0) : (!tf_res)
        -> !tf_res
      // CHECK-NEXT: "tf.Identity"
      // CHECK-SAME: {device = "/TPU:0"}
      %id1 = "tf.Identity"(%id0) : (!tf_res)
        -> !tf_res
      // CHECK-NEXT: "tf.Identity"
      // CHECK-SAME: {device = "/CPU:0"}
      %id2 = "tf.Identity"(%var_handle) : (!tf_res)

  // CHECK: input_sharding_configuration
  // CHECK-SAME: ["\08\01\1A\01\01\22\01\00"]
  // CHECK: output_sharding_configuration
  // CHECK-SAME: ["\08\01\1A\01\01\22\01\00"]
  func.return
}

// CHECK-LABEL: func @func_without_sharding
// CHECK-SAME: (%{{[a-z0-9]+}}: tensor<*xi32> {mhlo.sharding = "\08\01\1A\01\01\22\01\00"})
// CHECK-SAME: -> (tensor<*xi32> {mhlo.sharding = "\08\01\1A\01\01\22\01\00"})