input.getLoc(), input, rewriter.getDenseI64ArrayAttr(permutation));
}

// Defers the transpose of the left-hand side (LHS) to the right-hand side and
// the result of a binary operation. In detail, this rewrites the
// `op(transpose(%rhs), %lhs)` to `transpose(op(%rhs, transpose(%lhs)))`. The
// LHS transpose permutation must be a NCHW->NHWC permutation.
template <typename OpT>

	"os"
	"path/filepath"
	"sort"
	"strings"

	"cmd/internal/telemetry"
)

var (
	fset     = token.NewFileSet()
	exitCode = 0
)

var allowedRewrites = flag.String("r", "",
	"restrict the rewrites to this comma-separated list")

var forceRewrites = flag.String("force", "",
	"force these fixes to run even if the code looks updated")

var allowed, force map[string]bool

var (

#include <memory>

#include "llvm/ADT/StringRef.h"
#include "mlir/IR/BuiltinOps.h"  // from @llvm-project
#include "mlir/Pass/Pass.h"  // from @llvm-project

namespace mlir {
namespace TFTPU {

// Creates a pass that rewrites `tf_device.launch_func` on TPUs into TPU runtime
// ops.
std::unique_ptr<mlir::OperationPass<mlir::ModuleOp>> CreateTPURewritePass(
    llvm::StringRef module_name = llvm::StringRef());

#include "tensorflow/core/framework/types.h"
#include "tensorflow/core/lib/core/status.h"

namespace tensorflow {

// An OpKernel that compiles an op to an XLA computation and runs it. Unlike
// XlaLaunch this doesn't rely on any rewrites of the graphdef - it will run a
// vanilla TensorFlow op as long as the bridge supports it.
class XlaCompileOnDemandOp : public OpKernel {
 public:
  explicit XlaCompileOnDemandOp(OpKernelConstruction* ctx)

namespace mlir {
namespace TFDevice {

namespace {

#define GEN_PASS_DEF_CONVERTLAUNCHFUNCTOTFCALLPASS
#include "tensorflow/compiler/mlir/tensorflow/transforms/tf_passes.h.inc"

// Rewrites tf_device::LaunchFuncOp into TF::PartitionedCallOp.
struct ConvertLaunchFuncToTFCallPass
    : public impl::ConvertLaunchFuncToTFCallPassBase<
          ConvertLaunchFuncToTFCallPass> {
  void runOnOperation() override;
};

  pm.addPass(mlir::createInlinerPass());
  pm.addNestedPass<mlir::func::FuncOp>(mlir::createCanonicalizerPass());
  pm.addPass(mlir::quant::CreateCastBf16OpsToF32Pass());

  // Optimizes the graph via cleanups, merges, rewrites, constant folding,
  // and edge case handling where possible.
  pm.addNestedPass<mlir::func::FuncOp>(
      mlir::TF::CreateDropWhileShapeInvariantPass());
  pm.addNestedPass<mlir::func::FuncOp>(

            FindUserOfType<TFL::QuantizeOp>(op) != nullptr);
  }
};

// Rewrites `stablehlo.convolution` into fused `tfl.conv_2d`.
// If available, fuse bias and activation adjacent to `stablehlo.convolution`.
// This RewritePattern rewrites both the following into `tfl.conv_2d` op:
//
// StableHLO Quantizer output:
//   * input: per-tensor qi8

  let dependentDialects = ["tf_device::TensorFlowDeviceDialect"];
}

def XlaRewritePass : Pass<"tf-xla-rewrite", "mlir::ModuleOp"> {
  let summary = "Rewrites partition calls into Xla launch ops to make the attached function run on XLA.";

  let description = [{
    This pass rewrites `tf.PartitionedCall` and `tf.StatefulPartitionedCall`
    operations with `_xla_compile_device_type` attribute in a

// some generic types that are legal in MHLO. This pass legalizes TF types into
// types that are legal in MHLO. For example, TF::Qint8Type is converted to i8.
// Rewrites here should run before TF to MHLO op legalizations are run.

#include <memory>

#include "absl/log/log.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/Support/Casting.h"

include "mlir/Pass/PassBase.td"

def TPURewritePass : Pass<"tf-tpu-rewrite", "mlir::ModuleOp"> {
  let summary = "Rewrites a `tf_device.cluster_func` on TPUs into TPU runtime operations.";

  let description = [{
    This pass rewrites a `tf_device.cluster_func` operation into a sequence of `tf._TPUCompileMlir`
    and `tf.TPUExecute` operations. `tf._TPUCompileMlir` contains a MLIR module that is