// RUN: odml-to-stablehlo-opt -composite-lowering -verify-diagnostics %s | FileCheck %s

func.func @hardswish(%arg0: tensor<2xf32>) -> (tensor<*xf32>) {
  %0 = mhlo.composite "aten.hardswish.default" %arg0 {decomposition = @XlaCallModule_aten.hardswish.default.impl_0} : (tensor<2xf32>) -> tensor<2xf32>
  %1 = "tf.Identity"(%0) {device = ""} : (tensor<2xf32>) -> tensor<*xf32>
  %2 = "tf.Identity"(%1) {device = ""} : (tensor<*xf32>) -> tensor<*xf32>

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.
  }];

  // canonicalized to plain HLO constant if statically shaped. Add the
  // canonicalization pattern to pattern list to enable multi-hop lowering.
  chlo::ConstantLikeOp::getCanonicalizationPatterns(patterns, context);

  return ApplyPatterns(op, patterns, legalize_chlo);
}

// Performs the lowering to XLA dialect.
void LegalizeTF::runOnOperation() {
  auto op = getOperation();
  auto op_name = op->getName().getStringRef().str();

                          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 {

	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

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

    // IDE tries to acquire its bytecode via the index, however,
    // the index doesn't cover classfiles from compiler output,
    // so the lowering fails.
    //
    // To solve the problem, we need to find all delegated properties with inline accessors
    // reachable from files that will be compiled, and include files these inline accessors

  // 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>
}

// 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.

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.