// CHECK: tf.Atan2
    // expected-remark@+1 {{lowering requires bounded tensor operands}}
    %0 = "tf.Atan2"(%arg0, %arg0) : (tensor<*xf32>, tensor<*xf32>) -> tensor<*xf32>

    func.return %0 : tensor<*xf32>
  }

  // CHECK-LABEL: dynamic_operand
  func.func @dynamic_operand(%arg0: tensor<?xf32>) -> tensor<?xf32> {
    // CHECK: tf.Atan2
    // expected-remark@+1 {{lowering requires bounded tensor operands}}

	{name: "Arg", aux: "SymOff", symEffect: "Read", zeroWidth: true}, // argument to the function.  aux=GCNode of arg, off = offset in that arg.

	// Like Arg, these are generic ops that survive lowering. AuxInt is a register index, and the actual output register for each index is defined by the architecture.
	// AuxInt = integer argument index (not a register number). ABI-specified spill loc obtained from function

// is the responsibility of a clustering pass to combine all these individual
// operations constraints to form a valid cluster.
//
// Example: compilation using XLA (MHLO) lowering
//
//   %0 = "tf.Transpose"(%input, %perm)
//        : (tensor<?x?xf32>, tensor<2xi32>) -> tensor<?x?xf32>
//
//   XLAs `mhlo.transpose` operation requires permutation to be an attribute

    pm.addPass(
        mlir::odml::createPrintOpStatsPass(GetAcceptedStableHLODialects()));
  }

  if (failed(pm.run(tf_module))) {
    return tensorflow::errors::Aborted("Lowering to StableHLO failed.");
  }

  return absl::OkStatus();
}

tensorflow::Status RunConverter(const PassPipelineCLParser& pass_pipeline) {
  DialectRegistry registry;
  registerAllDialects(registry);

// faultOnLoad is true if a load to an address below minZeroPage will trigger a SIGSEGV.
var faultOnLoad = buildcfg.GOOS != "aix"

// nilcheckelim2 eliminates unnecessary nil checks.
// Runs after lowering and scheduling.
func nilcheckelim2(f *Func) {
	unnecessary := f.newSparseMap(f.NumValues()) // map from pointer that will be dereferenced to index of dereferencing value in b.Values[]
	defer f.retSparseMap(unnecessary)

  let description = [{
    The TensorFlow Device dialect.

    This dialect contains operations to describe/launch computations on devices.
    These operations do not map 1-1 to TensorFlow ops and requires a lowering
    pass later to transform them into Compile/Run op pairs, like XlaCompile and
    XlaRun.
}];

  let cppNamespace = "::mlir::tf_device";
}

        // Wrap these special attributes as a special TFR constant, so the SSA
        // value has a valid type to be used as TFR function argument. These
        // attributes are not expected to be manipulated by the lowering passes.
        if (mlir::isa<TypeAttr>(attribute) || mlir::isa<ArrayAttr>(attribute) ||
            mlir::isa<StringAttr>(attribute) ||
            mlir::isa<FlatSymbolRefAttr>(attribute)) {

          GetEinsumDimensionNumbersUnary(op.getEquation(), lhs)) {
    return rewriteToReduceSumAndTranspose(op, dnums_or.value(), rewriter);
  }

  return rewriter.notifyMatchFailure(op, "unsupported einsum lowering");
}

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

// Transform Einsum to other TF Ops for the supported variants.

// Copyright 2016 The Go Authors. All rights reserved.
// Use of this source code is governed by a BSD-style
// license that can be found in the LICENSE file.

// Lowering arithmetic
(Add(Ptr|64|32|16|8) ...) => (ADD ...)
(Add(64|32)F ...) => (FADD(D|S) ...)

(Sub(Ptr|64|32|16|8) ...) => (SUB ...)
(Sub(64|32)F ...) => (FSUB(D|S) ...)

(Mul64 ...) => (MUL  ...)
(Mul64uhilo ...) => (LoweredMuluhilo ...)
(Mul64uover ...) => (LoweredMuluover ...)

//	bb2            bb3
//
// where the intermediate blocks are mostly empty (with no side-effects);
// rewrite Phis in the postdominator as CondSelects.
func branchelim(f *Func) {
	// FIXME: add support for lowering CondSelects on more architectures
	switch f.Config.arch {
	case "arm64", "ppc64le", "ppc64", "amd64", "wasm", "loong64":
		// implemented
	default:
		return
	}