// CHECK-LABEL: func @tfAssertTrue
func.func @tfAssertTrue(%arg0: tensor<1x1x6x2xf32>) {
  %t = arith.constant dense<true> : tensor<i1>
  // CHECK-NOT: tf.Assert
  "tf.Assert"(%t, %arg0) {summarize = 3} : (tensor<i1>, tensor<1x1x6x2xf32>) -> ()
  func.return
}

// CHECK-LABEL: func @tfAssertFalse
func.func @tfAssertFalse(%arg0: tensor<1x1x6x2xf32>) {
  %f = arith.constant dense<false> : tensor<i1>
  // CHECK: tf.Assert

// CHECK:           })
// CHECK-SAME:        -> tensor<2x4x6x7xf32>
// CHECK:           %[[RESULT:.*]] = mhlo.divide %[[DIVIDEND]], %[[DIVISOR]] : tensor<2x4x6x7xf32>
// CHECK:           return %[[RESULT]] : tensor<2x4x6x7xf32>
// CHECK:         }
func.func @avgpool_same_padding(%arg0: tensor<2x12x21x7xf32>) -> tensor<2x4x6x7xf32> {

  // We cannot fuse this tfl.mul into the preceding conv op because %cst2 is not broadcast-compatible to %cst0.
  %1 = "tfl.mul"(%0, %cst2) {fused_activation_function = "RELU6"} : (tensor<1x4x4x2xf32>, tensor<4x2xf32>) -> tensor<1x4x4x2xf32>

  func.return %1 : tensor<1x4x4x2xf32>

// CHECK-NOT: tfl.gather_nd
// CHECK-NOT: tfl.gather

// -----

// Tests that a float `stablehlo.gather` is not converted to `tfl.gather_nd`.

func.func @gather_float(%arg0: tensor<3x4x2x2xf32>, %arg1: tensor<2x3x2xi64>) -> tensor<2x3x2x2xf32> {
  %0 = "stablehlo.gather"(%arg0, %arg1) {
    dimension_numbers = #stablehlo.gather<
      offset_dims = [2, 3],
      collapsed_slice_dims = [0, 1],

}

// -----

// CHECK-LABEL: topk_2
func.func @topk_2(%arg0: tensor<3x4x8xf32>) -> (tensor<3x4x2xf32>, tensor<3x4x2xi32>) {
  %0 = arith.constant dense<2> : tensor<i32>
  %1:2 = "tfl.topk_v2"(%arg0, %0) : (tensor<3x4x8xf32>, tensor<i32>) -> (tensor<3x4x2xf32>, tensor<3x4x2xi32>)
  func.return %1#0, %1#1: tensor<3x4x2xf32>, tensor<3x4x2xi32>
}

// -----

// CHECK-LABEL: topk_d