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 {

// license that can be found in the LICENSE file.

// Check to make sure we don't try to constant fold a divide by zero.
// This is a tricky test, as we need a value that's not recognized as 0
// until lowering (otherwise it gets handled in a different path).

package p

func f() {
	var i int
	var s string
	for i > 0 {
		_ = s[0]
		i++
	}

	var c chan int
	c <- 1 % i
}

# Target Aware Conversion (TAC)

Different hardwares have different capabilities and restrictions.

TAC is designed to leverage hardwares' capabilities to:

*   Perform device-specific optimizations (such as unsupported ops lowering,
    layout transformations, etc.)
*   Graph partitioning based on the hardware costs modeling.
*   It supports general import/export where you can hook your own

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

// Lowers some of the TensorFlow operations that can be represented using other
// TensorFlow operations.
struct LowerTF : public impl::TestTensorFlowLowerTFPassBase<LowerTF> {
  void runOnOperation() override {
    // Add lowering patterns to the list.
    RewritePatternSet patterns(&getContext());
    if (default_patterns_) {
      mlir::TF::PopulateLoweringTFPatterns(&getContext(), &patterns);
    }
    if (pre_hlo_patterns_) {

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

/// Creates a pass that converts constants followed by a qbarrier to a
/// constant whose value is quantized. This is typically one of the last
/// passes done when lowering to express actual quantized arithmetic in a
/// low level representation. Because it modifies the constant, it is
/// destructive and cannot be undone.
std::unique_ptr<OperationPass<func::FuncOp>> createConvertConstPass();

namespace mlir {
namespace tf_test {

// Returns test pass for variable freezing.
std::unique_ptr<OperationPass<ModuleOp>> CreateFreezeVariableTestPass();

// Test pass for applying TF->TF lowering patterns.
std::unique_ptr<OperationPass<func::FuncOp>> CreateTestTFLowerTFPass();

// Test passes for visitor util.
std::unique_ptr<OperationPass<func::FuncOp>> CreateTestVisitorUtilPass();