%1 = "tf.ReadVariableOp"(%0) {device = "/device:CPU:0"} : (tensor<!tf_type.resource<tensor<f32>>>) -> tensor<f32>
  %2 = "tf.XlaLaunch"(%arg0, %1) {_noinline = true, _xla_compile_device_type = "GPU", device = "/device:GPU:0", function = @xla_func_0, operandSegmentSizes = array<i32: 0, 2, 0>} : (tensor<7xf32>, tensor<f32>) -> tensor<10xf32>
  func.return %2 : tensor<10xf32>

// RUN: tf-opt %s --tf-print | FileCheck %s

module {
// Smoke test. We don't expect any modifications of the MLIR.

// CHECK-LABEL: foo
// CHECK: return
func.func @foo(%arg0: tensor<f32>, %arg1: tensor<f32>) -> tensor<f32> {
  return %arg0 : tensor<f32>
}

# function is successful.

# CHECK-LABEL: HloModule main
# CHECK:       ENTRY %main.{{[0-9]+}} ([[ARG0:.*]]: f32[]) -> (f32[]) {
# CHECK-NEXT:    %[[ARG0]] = f32[] parameter(0)
# CHECK-NEXT:    ROOT %tuple.{{[0-9]+}} = (f32[]) tuple(f32[] %[[ARG0]])

// CHECK-SAME: ([[ARG0:%.*]]: tensor<i32>, [[ARG1:%.*]]: tensor<f32>)
func.func @testIdentityN(%arg0: tensor<i32>, %arg1: tensor<f32>) {
  // CHECK-NOT:  "tf.IdentityN"
  // CHECK:      "tf_device.cluster"
  // CHECK-NEXT: tf_device.return [[ARG0]], [[ARG1]]
  %0:2 = "tf_device.cluster"() ({
    %1:2 = "tf.IdentityN"(%arg0, %arg1) : (tensor<i32>, tensor<f32>) -> (tensor<i32>, tensor<f32>)
    tf_device.return %1#0, %1#1 : tensor<i32>, tensor<f32>

    func.return %0 : tensor<?x19xf32>
  }
}

// CHECK-LABEL: HloModule main
// CHECK:       (arg_tuple.{{[0-9]+}}: (f32[10,17], f32[17,19])) -> (f32[10,19])

// NO_TUPLES-LABEL: HloModule main

// CHECK: %[[dq0:.*]] = "quantfork.dcast"(%[[q0]])
// CHECK-SAME: quant.uniform<i8:f32, 0.010039215461880554:-1>
// CHECK: %[[maxpool:.*]] = "tf.MaxPool"(%[[dq0]])
// CHECK: %[[q1:.*]] = "quantfork.qcast"(%[[maxpool]])
// CHECK-SAME: quant.uniform<i8:f32, 0.010039215461880554:-1>
// CHECK: %[[dq1:.*]] = "quantfork.dcast"(%[[q1]])
// CHECK-SAME: quant.uniform<i8:f32, 0.010039215461880554:-1>
// CHECK: %[[reshape:.*]] = "tf.Reshape"(%[[dq1]]

  func.func @f(%arg0: tensor<!tf_type.resource<tensor<f32>>> {tf_saved_model.bound_input = @v})
  attributes {tf_saved_model.exported_names = ["f"]} {
    // CHECK: "tf.ReadVariableOp"
    %0 = "tf.ReadVariableOp"(%arg0) {device = ""} : (tensor<!tf_type.resource<tensor<f32>>>) -> tensor<f32>
    func.return
  }

// RUN: flatbuffer_translate -mlir-to-tflite-flatbuffer %s -o - | flatbuffer_to_string - | FileCheck %s

func.func @main(%arg0: tensor<*x!quant.uniform<u16:f32, 2.0:37>>) -> tensor<*x!quant.uniform<u16:f32, 2.0:37>> {
// CHECK:     {
// CHECK-NEXT:  version: 3,
// CHECK-NEXT:  operator_codes: [ ],
// CHECK-NEXT:  subgraphs: [ {
// CHECK-NEXT:    tensors: [ {
// CHECK-NEXT:      shape: [  ],
// CHECK-NEXT:      type: UINT16,
// CHECK-NEXT:      buffer: 1,

// CHECK: %[[CST:.+]] = stablehlo.constant() <{value = dense<127> : tensor<2x3xi8>}> : () -> tensor<2x3x!quant.uniform<i8<-127:127>:f32, 0.0023622048182750312>>
// CHECK: %[[CALL:.+]] = call @quantized_dot_general_fn(%[[ARG0]], %[[CST]]) {_quantization_method = "weight_only_ptq { }"} : (tensor<1x2xf32>, tensor<2x3x!quant.uniform<i8<-127:127>:f32, 0.0023622048182750312>>) -> tensor<1x3xf32>
// CHECK: return %[[CALL]]

// CHECK: quantized_dot_general_fn

  // CHECK: %[[cst:.*]] = stablehlo.constant
  // CHECK: %[[q1:.*]] = "quantfork.qcast"(%[[cst]])
  // CHECK-SAME: quant.uniform<i8:f32, 0.0040316890267764818:127>
  // CHECK: %[[dq1:.*]] = "quantfork.dcast"(%[[q1]])
  // CHECK-SAME: quant.uniform<i8:f32, 0.0040316890267764818:127>
  %cst = stablehlo.constant dense<[[-0.960978984, -0.390246302], [-0.790828585, -0.601039409], [-1.0280807, -1.02731466]]> : tensor<3x2xf32>