include "mlir/Pass/PassBase.td"

// TF dialect passes.

def TensorflowGPUFusion : Pass<"tf-gpu-op-fusion", "mlir::func::FuncOp"> {
  let summary = "Fusion optimization for GPU targets";
  let description = [{
    This pass is performing fusion specific to GPU targets. This is an ad-hoc
    pass for now, but should be integrated with some notion of "target" in the
    MLIR pipeline in the future.

        /*value=*/constant_op.getValue());
  }
}

// Creates a new `tfl.qconst` op for the bias. The bias values are 0s, because
// this bias a dummy bias (note that bias fusion is not considered for this
// transformation). The quantization scale for the bias is input scale *
// filter scale. `filter_const_op` is used to retrieve the filter scales and
// the size of the bias constant.

    if (llvm::isa<AddOp, SubOp>(binary_op)) {
      auto padding = fc_op->template getAttrOfType<StringAttr>("padding");
      if (padding && padding.getValue() != "VALID") return failure();

      // The fusion of add/sub is actually applying the following
      // transformation:
      // w * (x + c) + b => w * x + (w * c + b)
      // so we have to update the bias.
      if (llvm::isa<SubOp>(binary_op)) cst_value.changeSign();

// -----

// Tests that a fused per-channel quantized `stablehlo.dot_general` is properly
// lowered to fused `tfl.fully_connected`.
// TODO: b/309896242 - Add more support for dynamic bias fusion cases.

func.func @dot_general_with_bias_same_shape_srq(%arg0: tensor<1x1024x!quant.uniform<i8:f32, 1.000000e+0:0>>) -> (tensor<1x3x!quant.uniform<i8:f32, 3.000000e+00:-127>>) {

  auto operands = adaptor.getOperands();
  // TODO(b/142478136): Handle fused ops.
  if (getFusedActivationFunction() != "NONE") return {};

  // This function is performance critical for op fusion patterns, e.g.
  // FuseBinaryOpToPrecedingAffine and FuseMulOrDivWithConv2dOrDepthwiseConv2d.
  // So a few specializations are provided to evaluate the math operation
  // more efficiently.

}

// fusedJump determines whether p can be fused with a subsequent conditional jump instruction.
// If it can, we return true followed by the total size of the fused jump. If it can't, we return false.
// Macro fusion rules are derived from the Intel Optimization Manual (April 2019) section 3.4.2.2.
func fusedJump(p *obj.Prog) (bool, uint8) {
	var fusedSize uint8

from the value obtained by executing and rounding the instructions individually.
An explicit floating-point type <a href="#Conversions">conversion</a> rounds to
the precision of the target type, preventing fusion that would discard that rounding.
</p>

<p>
For instance, some architectures provide a "fused multiply and add" (FMA) instruction

// fans : 2014-11-07 ZDNS International Limited
fans

// farm : 2013-11-07 Binky Moon, LLC
farm

// farmers : 2015-07-09 Farmers Insurance Exchange
farmers

// fashion : 2014-07-03 Registry Services, LLC
fashion

// fast : 2014-12-18 Amazon Registry Services, Inc.
fast

// fedex : 2015-08-06 Federal Express Corporation
fedex

// feedback : 2013-12-19 Top Level Spectrum, Inc.

   * ListenableFuture instances
   */

  /**
   * A future in {@link TestFutureBatch} that also has a name for debugging purposes and a {@code
   * finisher}, a task that will complete the future in some fashion when it is called, allowing for
   * testing both before and after the completion of the future.
   */
  @J2ktIncompatible
  @GwtIncompatible // used only in GwtIncompatible tests

   * ListenableFuture instances
   */

  /**
   * A future in {@link TestFutureBatch} that also has a name for debugging purposes and a {@code
   * finisher}, a task that will complete the future in some fashion when it is called, allowing for
   * testing both before and after the completion of the future.
   */
  @J2ktIncompatible
  @GwtIncompatible // used only in GwtIncompatible tests