" a verification pass to ensure invariants are true."}];

  let options = [
    Option<"legalize_chlo_", "legalize-chlo", "bool", /*default=*/"true",
        "Legalizes intermediate chlo ops to hlo">
    ];

  let constructor = "mlir::mhlo::CreateVerifyTFXLALegalizationPass()";
}

def TFXLADeviceSpecificTransforms : Pass<"tfxla-device-specific-transforms",

#include "tensorflow/compiler/mlir/tf2xla/transforms/passes.h"
#include "xla/mlir_hlo/mhlo/IR/hlo_ops.h"
#include "xla/mlir_hlo/mhlo/IR/register.h"
#include "xla/mlir_hlo/mhlo/transforms/passes.h"
#include "xla/mlir_hlo/mhlo/transforms/rewriters.h"
#include "xla/mlir_hlo/mhlo/utils/type_conversion.h"

namespace mlir {
namespace odml {

class TFToMhloPass

  // CHECK-DAG: %[[RES:.*]] = mhlo.constant dense<{{\[\[\[\[}}1, 4, 9, 16], [5, 12, 21, 32]]]]> : tensor<1x1x2x4xi32>
  %cst0 = mhlo.constant dense<[1, 2, 3, 4]> : tensor<4xi32>
  %cst1 = mhlo.constant dense<[[[[1, 2, 3, 4], [5, 6, 7, 8]]]]> : tensor<1x1x2x4xi32>
  %0 = "mhlo.broadcast_in_dim"(%cst0) <{broadcast_dimensions = dense<3> : tensor<1xi64>}> : (tensor<4xi32>) -> tensor<1x1x2x4xi32>
  %1 = mhlo.multiply %0, %cst1 : tensor<1x1x2x4xi32>

  // CHECK-DAG: %[[FILTER:.+]] = mhlo.constant dense<{{\[\[\[\[}}1.000000e+00, 2.000000e+00], [3.000000e+00, 4.000000e+00], [5.000000e+00, 6.000000e+00]]]]> : tensor<1x1x3x2xf32>
  // CHECK-DAG: %[[CST_0:.+]] = mhlo.constant dense<[1.000000e-01, 2.000000e-01]> : tensor<2xf32>
  // CHECK-DAG: %[[CST_1:.+]] = mhlo.constant dense<[3.000000e-01, 4.000000e-01]> : tensor<2xf32>

        return
      }
    ```

    This might be legalized like this:

    ```mlir
      func @send_to_host(%arg0: tensor<i32>) {
        %0 = "mhlo.create_token"() : () -> !mhlo.token
        %1 = "mhlo.send"(%arg0, %0) {...}} : (tensor<i32>, !mhlo.token) -> !mhlo.token
        return
      }
    '''
  }];

func.func @mhlo_custom_call_test__dont_legalize_dict_backend_config(%arg0: tensor<1x4xf32>) -> tensor<1x8xf32> {
  %0 = mhlo.custom_call @custom_call.my_custom_op(%arg0) {
    api_version = 4 : i32,
    backend_config = {foo = "bar"}
  } : (tensor<1x4xf32>) -> (tensor<1x8xf32>)
  func.return %0 : tensor<1x8xf32>

  //       CHECK: %0 = mhlo.custom_call @custom_call.my_custom_op(%arg0) {
  //  CHECK-SAME:   api_version = 4 : i32,

// like tf.ResizeBilinear or tf.ResizeNearestNeighbor to mhlo.custom_call ops.
void AddTFToStablehloPasses(OpPassManager& pm, bool skip_resize,
                            bool smuggle_disallowed_ops);

// This function is a common entry point for all graph optimizations that are
// not specific to any hardware. It legalizes SHLO->MHLO, does MHLO->MHLO
// optimizations by calling `AddMhloOptimizationPasses` internally, and

// CHECK:           [[LHSRE:%.*]] = mhlo.real [[LHS]]
// CHECK:           [[LHSIM:%.*]] = mhlo.imag [[LHS]]
// CHECK:           [[LHSIMNEG:%.*]] = mhlo.negate [[LHSIM]]
// CHECK:           [[LHSCONJ:%.*]] = mhlo.complex [[LHSRE]], [[LHSIMNEG]]
// CHECK:           [[RHSRE:%.*]] = mhlo.real [[RHS]]
// CHECK:           [[RHSIM:%.*]] = mhlo.imag [[RHS]]
// CHECK:           [[RHSIMNEG:%.*]] = mhlo.negate [[RHSIM]]

#include "tensorflow/compiler/mlir/quantization/common/attrs_and_constraints.h"
#include "xla/mlir_hlo/mhlo/IR/hlo_ops.h"

namespace mlir {
namespace odml {

class FuseMhloMulAndConvolutionPattern : public OpRewritePattern<mhlo::MulOp> {
 public:
  using OpRewritePattern<mhlo::MulOp>::OpRewritePattern;

  LogicalResult matchAndRewrite(mhlo::MulOp mul_op,
                                PatternRewriter &rewriter) const override {

// 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>
  // CHECK: %1 = "mhlo.broadcast_in_dim"(%0) <{broadcast_dimensions = dense<> : tensor<0xi64>}> : (tensor<f32>) -> tensor<1x750xf32>