}

// CHECK-LABEL: @StridedSliceEllipsisMaskBefore
func.func @StridedSliceEllipsisMaskBefore(%arg0: tensor<21x15x7xf32>) -> tensor<21x15x2xf32> {
  %cst = arith.constant dense<0> : tensor<2xi32>
  %cst_0 = arith.constant dense<1> : tensor<2xi32>

  %1 = "tfl.add"(%0, %cst) {fused_activation_function = "NONE"} : (tensor<256x8x7x16xf16>, tensor<f16>) -> tensor<256x8x7x16xf16>
  func.return %1 : tensor<256x8x7x16xf16>

  %arg19: tensor<1x640xf32>,
  %arg20: tensor<1x2048xf32>
) -> tensor<1x640xf32> {
  %cst = "tfl.no_value"() {value = unit} : () -> none
  %zero = "tfl.pseudo_const"() {value = dense<0.0> : tensor<640xf32>} : () -> tensor<640xf32>
  %0 = "tfl.lstm"(%arg0, %arg1, %arg2, %arg3, %arg4, %arg5, %arg6, %arg7, %arg8, %cst, %cst, %cst, %arg9, %arg10, %arg11, %arg12, %arg13, %zero, %arg19, %arg20, %arg15, %arg16, %arg17, %arg18) ({}) {

//
// cst = value that is > 0
// rem = remainder(arg0, cst)
// for i in 0 to len(arg1):
//    if (rem[i] < 0 && rem[i] != 0)
//       rem[i] += cst
// return rem
def : Pat<(MHLO_SelectOp
            (MHLO_AndOp
              (MHLO_CompareOp:$rltz
                (MHLO_RemOp:$rem $arg, (MHLO_ConstantOp $cst)),
                (MHLO_ConstantOp $cst1),

func.func @dot_int4(%arg0: tensor<?x3xf32>) -> tensor<?x2xf32> {
  // 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>

    "tf.LookupTableImportV2"(%0, %cst, %cst_0) {_has_manual_control_dependencies = true, device = ""} : (tensor<!tf_type.resource>, tensor<3x!tf_type.string>, tensor<3xi64>) -> ()
    %1 = "tf.HashTableV2"() {container = "", device = "", key_dtype = !tf_type.string, shared_name = "hash_table_1", use_node_name_sharing = false, value_dtype = i64} : () -> tensor<!tf_type.resource>

func.func @noinline0() -> (tensor<0xf32>) {
  %cst = "tf.Const"() {value = dense<1.0> : tensor<0xf32>} : () -> tensor<0xf32>
  return %cst : tensor<0xf32>
}

// -----

// Tests that the function not listed in the option `noinline-functions`
// is not marked tf._noinline = true.

// CHECK-LABEL: @inline
// CHECK-NOT: tf._noinline
func.func @inline() -> (tensor<0xf32>) {

  func.func @conv2d_unmatching_unit(%arg0: tensor<1x3x4x3xf32>) -> (tensor<1x3x2x2xf32>) {
    %cst = "tf.Const"() {device = "", value = dense<1.0> : tensor<2x3x3x2xf32>} : () -> tensor<2x3x3x2xf32>
    %0 = "tf.Cast"(%arg0) {Truncate = false, device = ""} : (tensor<1x3x4x3xf32>) -> tensor<1x3x4x3xf32>

    return %8 : tensor<*xf32>
  }

// CHECK-LABEL: conv_with_bias_and_relu
// CHECK-DAG: %[[cst:.*]] = arith.constant dense<[7.11401462, 7.05456924]> : tensor<2xf32>
// CHECK-DAG: %[[cst_1:.*]] = arith.constant dense<{{.*}}> : tensor<2x3x3x2xf32>

// CHECK: %[[q0:.*]] = "quantfork.qcast"(%[[cst]]) {volatile}
// CHECK-SAME: tensor<2x!quant.uniform<i32:f32, 0.044169864606680966>>

// CHECK-LABEL: @unfold_splat_constant_float
func.func @unfold_splat_constant_float() -> tensor<1x750xf32> {
  %cst = mhlo.constant dense<7.680000e+02> : tensor<1x750xf32>
  func.return %cst : tensor<1x750xf32>

  // CHECK-DAG: %0 = mhlo.constant dense<7.680000e+02> : tensor<f32>