// CHECK: %[[VAL_1:.*]] = "tf.Identity"(%[[VAL_0]]) : (tensor<*xf32>) -> tensor<*xf32>
  // CHECK: return %[[VAL_1]]
  %0 = "tf.Conv2D"(%arg2, %arg1) {data_format = "NHWC", dilations = [1, 1, 1, 1], explicit_paddings = [], padding = "SAME", strides = [1, 1, 1, 1], use_cudnn_on_gpu = true} : (tensor<8x32x32x3xf32>, tensor<1x1x3x128xf32>) -> tensor<*xf32>

  %1 = "tf.Transpose"(%arg0, %0) : (tensor<1x3x32x32xf32>, tensor<4xi32>) -> tensor<1x32x32x3xf32>

  // Compute in NHWC
  %2 = "tf.Conv2D"(%1, %arg1)
    {
      data_format = "NHWC",
      padding = "SAME",
      strides = [1, 1, 1, 1],
      dilations = [1, 1, 1, 1]
    } : (tensor<1x32x32x3xf32>, tensor<1x1x3x8xf32>) -> tensor<1x32x32x8xf32>

  // Convert result back: NHWC -> NCHW

         explicit_paddings = [1, 2, 3, 4, 5, 6, 7, 8],
         padding = "EXPLICIT",
         strides = [5, 6, 7, 8]
       } : (tensor<4xi32>, tensor<1x1x3x8xf32>, tensor<1x32x32x8xf32>)
         -> tensor<1x32x32x3xf32>

  func.return %0 : tensor<1x32x32x3xf32>
}

// CHECK-LABEL: func @transposeFusedBatchNormV3
func.func @transposeFusedBatchNormV3(
  %arg0: tensor<1x28x28x64xf32>,
  %arg1: tensor<64xf32>

  module attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 268 : i32}} {
    func.func @main() -> tensor<3x32x32x3xf32> {
      %0 = mhlo.constant dense<2.550000e+02> : tensor<3x32x32x3xf32>
      return %0 : tensor<3x32x32x3xf32>
    }
  })";
  CreateModule(kMlirModuleStr);

  auto result = Run();

  EXPECT_TRUE(result.failed());
}

}  // namespace

  %cst_0 = arith.constant dense<1.000000e+00> : tensor<f32>
  %cst_1 = arith.constant dense<0.000000e+00> : tensor<8xf32>
  %cst_2 = arith.constant dense<0.000000e+00> : tensor<8x3x3x3xf32>
  %0 = "tfl.sub"(%arg0, %cst_0) {fused_activation_function = "NONE"} : (tensor<3x15x14x3xf32>, tensor<f32>) -> tensor<3x15x14x3xf32>

// -----

func.func @testConv2DNoBiasCPU(%arg0: tensor<128x32x32x3xf32>, %arg1: tensor<64x3x3x3xf32>, %arg2: none) -> tensor<128x32x32x64xf32> {
  // CHECK: tac.cost = 0x4DD80000
  %0 = "tfl.conv_2d"(%arg0, %arg1, %arg2) {dilation_h_factor = 1 : i32, dilation_w_factor = 1 : i32, padding = "SAME", stride_h = 1 : i32, stride_w = 1 : i32, fused_activation_function = "RELU6", tac.device = "CPU"} : (tensor<128x32x32x3xf32>, tensor<64x3x3x3xf32>, none) -> tensor<128x32x32x64xf32>

  func.return %2 : tensor<1x32x32x16x!quant.uniform<i8:f32, 1.0>>
}

// -----

// CHECK-LABEL: @test_conv2d_replace_float
func.func @test_conv2d_replace_float(%arg0: tensor<1x32x32x8xf32>) -> tensor<1x32x32x16xf32> {

  func.return %0 : tensor<256x32x32x16xf32>
}

func.func @testConv2DDynamicShape(tensor<?x32x32x3xf32>, tensor<16x3x3x3xf32>, tensor<16xf32>) -> tensor<?x32x32x16xf32> {
^bb0(%arg0: tensor<?x32x32x3xf32>, %arg1: tensor<16x3x3x3xf32>, %arg2: tensor<16xf32>):
  // CHECK: _arithmetic_count = -1 : i64

// -----

func.func @ensureBiasForConv2d(%arg0: tensor<128x32x32x3xf32>, %arg1: tensor<32x1x1x3xf32>) -> tensor<128x32x32x32xf32> {
  %cst = "tfl.no_value"() {value = unit} : () -> none
  %0 = "tfl.conv_2d"(%arg0, %arg1, %cst) {dilation_h_factor = 1 : i32, dilation_w_factor = 1 : i32, fused_activation_function = "NONE", padding = "VALID", stride_h = 1 : i32, stride_w = 1 : i32} : (tensor<128x32x32x3xf32>, tensor<32x1x1x3xf32>, none) -> tensor<128x32x32x32xf32>

  // CHECK-SAME: dilations = [1, 3, 4, 2]
  // CHECK-SAME: explicit_paddings = [1, 2, 5, 6, 7, 8, 3, 4]
  // CHECK-SAME: padding = "EXPLICIT"
  // CHECK-SAME: strides = [5, 7, 8, 6]
  // CHECK-SAME: (tensor<1x32x32x3xf32>, tensor<1x1x3x8xf32>) -> tensor<1x7x6x8xf32>

  // CHECK: %[[RES_PERM:.*]] = "tf.Const"() <{value = dense<[0, 3, 1, 2]> : tensor<4xi64>}>
  // CHECK: %[[RES_TRANSPOSE:[0-9]*]] = "tf.Transpose"(%[[CONV2D]], %[[RES_PERM]])