def InputLoweringMetricsPass : Pass<"input-lowering-metrics-pass", "mlir::func::FuncOp"> {

  let summary = "Collects various metrics about the input to the lowering "
                "portion of the bridge. This is a logical no-op.";

  let description = [{
    Gathers metrics about the input MLIR to Phase 2 of the TFXLA Bridge, which
    does a strict semantic lowering from Tensorflow ops to XLA HLO.
  }];

}

// BatchMatMulV2 has native as well as fallback lowering patterns available.
// The fallback pattern uses dot_general without broadcast on operands and then
// transposes the output which is faster. However, the fallback pattern doesn't
// support dynamic shaped operands like the native lowering. Verify that
// fallback lowering is preferred for static shaped operands when available.

// CHECK-LABEL: batchmatmulv2

                          tensorflow::ConvertTensor(tensor, &builder));
      builder.setInsertionPoint(
          &func_op.getFunctionBody().getBlocks().front().front());
      // Use mhlo.Constant when it is consumed by the lowering passes since they
      // can't lower tf.Const.
      Value cst;
      if (use_mhlo_const) {
        cst = builder.create<mhlo::ConstantOp>(func_op->getLoc(), attrs);
      } else {

bool IsOpAllowedTf2xlaFallback(const TypeID& type_id);

// Whether this type is Preferred to use a TF2XLA fallback kernel when using
// the MLIR bridge. If this is true, then the TF2XLA fallback kernel will be
// used over the MLIR lowering.
bool IsOpAllowedTf2xlaPreferred(const TypeID& type_id);

}  // namespace mhlo
}  // namespace mlir

	Types          Types
	lowerBlock     blockRewriter  // block lowering function, first round
	lowerValue     valueRewriter  // value lowering function, first round
	lateLowerBlock blockRewriter  // block lowering function that needs to be run after the first round; only used on some architectures
	lateLowerValue valueRewriter  // value lowering function that needs to be run after the first round; only used on some architectures

struct DecomposeResourceOpsPass
    : public impl::DecomposeResourceOpsPassBase<DecomposeResourceOpsPass> {
  void runOnOperation() override {
    // Add lowering patterns to the list.
    RewritePatternSet patterns(&getContext());
    TF::PopulateDecomposeResourceOpsPatterns(&getContext(), &patterns);

    if (failed(applyPatternsAndFoldGreedily(getOperation(),

  // CHECK: mlrt.await_handle [[handle_1]]
  mlrt.await_handle %handle_0
  mlrt.await_handle %handle_1

  // CHECK: return [[result]]
  return %result : tensor<i32>
}

// -----

// Test lowering tf.If

func.func @then(%x: tensor<i32>, %y: tensor<i32>) -> tensor<i32> {
  return %x: tensor<i32>
}

func.func @else(%x: tensor<i32>, %y: tensor<i32>) -> tensor<i32> {
  return %y: tensor<i32>
}

using ::tensorflow::thread::ThreadPoolInterface;

using ::tensorflow::tfrt_stub::FallbackTensor;

// -------------------------------------------------------------------------- //
// Run function via the TF->TFRT fallback lowering.
// -------------------------------------------------------------------------- //

namespace {
// Thread pool for running `intra-op` tasks scheduled by the fallback kernels.

// Creates an instance of the VerifyQuantLegalization pass, which verifies all
// quant ops and types are lowered.
std::unique_ptr<OperationPass<func::FuncOp>>
CreateVerifyQuantLegalizationPass();

// Add all passes for lowering TF quant ops and types to MHLO int.
void AddQuantizationLoweringPasses(mlir::OpPassManager &pm);

// Creates an instance of OptimizeIntGraphPass, which optimizes the int graph
// lowered from the quantized graph.

(Neg64F ...) => (FNEG ...)

(Com64 ...) => (NOT ...)
(Com(32|16|8) ...) => (NOTW ...)
(NOT x) => (XOR (MOVDconst [-1]) x)
(NOTW x) => (XORWconst [-1] x)

// Lowering boolean ops
(AndB ...) => (ANDW ...)
(OrB ...) => (ORW ...)
(Not x) => (XORWconst [1] x)

// Lowering pointer arithmetic
(OffPtr [off] ptr:(SP)) => (MOVDaddr [int32(off)] ptr)
(OffPtr [off] ptr) && is32Bit(off) => (ADDconst [int32(off)] ptr)