// CHECK-SAME: , asymmetric_quantize_inputs = true
// CHECK: return %[[mm:.*]]

// PerTensor: %[[w:.*]] = "tfl.pseudo_qconst"() <{qtype = tensor<512x12x!quant.uniform<i8<-127:127>:f32, 1.000000e+00>>,
// PerTensor: %[[mm:.*]] = "tfl.batch_matmul"(%arg0, %[[w]]) <{adj_x = false, adj_y = false
// PerTensor-SAME: , asymmetric_quantize_inputs = true
// PerTensor: return %[[mm:.*]]
}

  // FB-SAME:  -8.0, -9.0, -10.0, -11.0, -12.0, -13.0, -14.0, -15.0, -16.0,
  // FB-SAME:  -17.0, -18.0, -19.0, -20.0, -21.0, -22.0, -23.0, -24.0, -25.0,
  // FB-SAME:  -26.0, -27.0, -28.0, -29.0, -30.0, -31.0, -32.0, -33.0, -34.0,
  // FB-SAME:  -35.0, -36.0, -37.0, -38.0, -39.0 ],
  // FB-NEXT:        max: [ 0.0, 1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0,

  func.return %3 : tensor<4x8xf32>
}

// CHECK-LABEL:   func @select(
// CHECK-SAME:                 %[[VAL_0:.*]]: tensor<2xi1>,
// CHECK-SAME:                 %[[VAL_1:.*]]: tensor<2xi32>,
// CHECK-SAME:                 %[[VAL_2:.*]]: tensor<2xi32>) -> tensor<2xi32> {

  func.return
}

// CHECK:      tf_device.replicate
// CHECK-SAME: [%[[ARG_0]], %[[ARG_0]]] as %{{[a-z0-9]*}}
// CHECK-SAME: [%[[ARG_2]], %[[ARG_2]]] as %{{[a-z0-9]*}}
// CHECK-SAME: [%[[ARG_4]], %[[ARG_4]]] as %{{[a-z0-9]*}}
// CHECK-SAME: %[[ARG_1]] as %{{[a-z0-9]*}}
// CHECK-SAME: %[[ARG_3]] as %{{[a-z0-9]*}}
// CHECK-SAME: %[[ARG_5]] as %{{[a-z0-9]*}}

// CHECK: return %[[CALL]]

// CHECK: quantized_dot_general_fn
// CHECK-SAME: (%[[ARG1:.+]]: tensor<1x2xf32>,  %[[ARG2:.+]]: tensor<2x3x!quant.uniform<i8<-127:127>:f32, 0.0023622048182750312>>) -> tensor<1x3xf32>
// CHECK: %[[DOT:.+]] = stablehlo.dot_general %[[ARG1]], %[[ARG2]]
// CHECK-SAME: (tensor<1x2xf32>, tensor<2x3x!quant.uniform<i8<-127:127>:f32, 0.0023622048182750312>>) -> tensor<1x3xf32>

  // 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]

  // 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]}

  %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: func private @[[CLUSTER]]
// CHECK-SAME: (%[[CLUSTER_ARG_0:[a-z0-9]*]]: tensor<?xi32>) -> tensor<?xi32>
// CHECK: %[[B_OUTPUT:[0-9]*]] = "tf.B"(%[[CLUSTER_ARG_0]])
// CHECK: return %[[B_OUTPUT]]

// -----

// Tests that multiple `tf_device.cluster` that depend on each other are
// correctly handled.

// CHECK-LABEL: func @multiple_clusters
// CHECK-SAME: (%[[ARG_0:[a-z0-9]*]]: tensor<?xi32>)

  func.return %1 : tensor<4x128x1xf32>
  // CHECK-NEXT: %[[CONST_WEIGHT:.*]] = arith.constant
  // CHECK-SAME: [1.000000e+00, 2.000000e+00]
  // CHECK-SAME: tensor<1x2xf32>
  // CHECK: %[[FC_RES:.*]] = "tfl.fully_connected"(%arg0, %[[CONST_WEIGHT]]
  // CHECK-SAME: <{fused_activation_function = "NONE", keep_num_dims = true, weights_format = "DEFAULT"}> : (tensor<4x128x2xf32>, tensor<1x2xf32>, none) -> tensor<4x128x1xf32>