// CHECK: %[[x_backprop:.*]] = mhlo.convert %[[mul3]] : tensor<8x8x8x8xf32>
    // CHECK: return %[[x_backprop]] : tensor<8x8x8x8xf32>

  // CHECK-NEXT: %[[x_backprop:.*]] = mhlo.convert %[[mul3]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: return %[[x_backprop]] : tensor<8x8x8x8xf32>

          {"BN",
           {"FusedBatchNorm", "FusedBatchNormV2", "FusedBatchNormV3",
            "_FusedBatchNormEx", "FusedBatchNormGrad", "FusedBatchNormGradV2",
            "FusedBatchNormGradV3"}},
          {"Conv", {"_FusedConv2D"}},
          {"SORT", {"TopKV2"}},  // XLA version much faster then TF version.
          {"MISC",
           // clang-format off

  );

  TF_DerivedOperandTypeAttr T = TF_DerivedOperandTypeAttr<0>;
  TF_DerivedOperandTypeAttr U = TF_DerivedOperandTypeAttr<3>;
}

def TF_FusedBatchNormGradV3Op : TF_Op<"FusedBatchNormGradV3", [Pure, TF_LayoutSensitiveInterface]> {
  let summary = "Gradient for batch normalization.";

  let description = [{
Note that the size of 4D Tensors are defined by either "NHWC" or "NCHW".