module {

  // Note: following functions won't handle per-channel quantization for now.
  func.func private @internal_quantize_i8(%input : tensor<*xf32>, %scale : tensor<*xf32>, %zp : tensor<*xi32>) -> tensor<*xi8> {
    // Uses tf.floor(x + 0.5) instead of tf.round(x) since tf.round generates
    // a very expensive pattern.
    %round_cst = "tf.Const"() {value = dense<0.5> : tensor<f32>} : () -> tensor<f32>

    @CheckForNull
    private AvlNode<E> floor(Comparator<? super E> comparator, @ParametricNullness E e) {
      int cmp = comparator.compare(e, getElement());
      if (cmp > 0) {
        return (right == null) ? this : MoreObjects.firstNonNull(right.floor(comparator, e), this);
      } else if (cmp == 0) {
        return this;
      } else {
        return (left == null) ? null : left.floor(comparator, e);
      }
    }

def LowerTruncateDivOpOnIntTensors : Pat<
  (TF_TruncateDivOp TF_IntTensor:$lhs, $rhs),
  (TF_DivOp $lhs, $rhs)>;

// Note: truncation could also be implemented as sign(x) * floor(abs(x)) or
//       (-1 & x) || floor(abs(x)), based on performance benchmarks.
def LowerTruncateDivOpOnFloatTensors : Pat<
  (TF_TruncateDivOp TF_FloatTensor:$lhs, $rhs),
  (TF_SelectV2Op
    (TF_LessOp
      (TF_DivOp:$div $lhs, $rhs),

    assertThat(set.floor("f")).isNull();
  }

  public void testFloor_elementPresent() {
    ImmutableSortedSet<String> set =
        ImmutableSortedSet.copyOf(new String[] {"e", "a", "e", "f", "b", "i", "d", "a", "c", "k"});
    assertThat(set.floor("f")).isEqualTo("f");
    assertThat(set.floor("j")).isEqualTo("i");
    assertThat(set.floor("q")).isEqualTo("k");
  }

    @CheckForNull
    private AvlNode<E> floor(Comparator<? super E> comparator, @ParametricNullness E e) {
      int cmp = comparator.compare(e, getElement());
      if (cmp > 0) {
        return (right == null) ? this : MoreObjects.firstNonNull(right.floor(comparator, e), this);
      } else if (cmp == 0) {
        return this;
      } else {
        return (left == null) ? null : left.floor(comparator, e);
      }
    }

		{name: "F32Ceil", asm: "F32Ceil", argLength: 1, reg: fp32_11, typ: "Float32"},         // ceil(arg0)
		{name: "F32Floor", asm: "F32Floor", argLength: 1, reg: fp32_11, typ: "Float32"},       // floor(arg0)
		{name: "F32Nearest", asm: "F32Nearest", argLength: 1, reg: fp32_11, typ: "Float32"},   // round(arg0)
		{name: "F32Abs", asm: "F32Abs", argLength: 1, reg: fp32_11, typ: "Float32"},           // abs(arg0)

	// In general (x*n)/2⁶⁴ = k for x*n in [k*2⁶⁴,(k+1)*2⁶⁴).
	// There are either floor(2⁶⁴/n) or ceil(2⁶⁴/n) possible products
	// in that range, depending on k.
	// But suppose we reject the sample and try again when
	// x*n is in [k*2⁶⁴, k*2⁶⁴+(2⁶⁴%n)), meaning rejecting fewer than n possible
	// outcomes out of the 2⁶⁴.
	// Now there are exactly floor(2⁶⁴/n) possible ways to produce
	// each output value k, so we've restored uniformity.

      throw new ArithmeticException("input is infinite or NaN");
    }
    switch (mode) {
      case UNNECESSARY:
        checkRoundingUnnecessary(isMathematicalInteger(x));
        return x;

      case FLOOR:
        if (x >= 0.0 || isMathematicalInteger(x)) {
          return x;
        } else {
          return (long) x - 1;
        }

      case CEILING:
        if (x <= 0.0 || isMathematicalInteger(x)) {

      throw new ArithmeticException("input is infinite or NaN");
    }
    switch (mode) {
      case UNNECESSARY:
        checkRoundingUnnecessary(isMathematicalInteger(x));
        return x;

      case FLOOR:
        if (x >= 0.0 || isMathematicalInteger(x)) {
          return x;
        } else {
          return (long) x - 1;
        }

      case CEILING:
        if (x <= 0.0 || isMathematicalInteger(x)) {

  %9 = mhlo.multiply %7, %8 : tensor<4xf32>
  %10 = "mhlo.broadcast_in_dim"(%2) <{broadcast_dimensions = dense<> : tensor<0xi64>}> : (tensor<f32>) -> tensor<4xf32>
  %11 = mhlo.divide %9, %10 : tensor<4xf32>
  %12 = mhlo.floor %11 : tensor<4xf32>
  %13 = mhlo.convert %12 : (tensor<4xf32>) -> tensor<4xi32>
  %14 = "mhlo.broadcast_in_dim"(%1) <{broadcast_dimensions = dense<> : tensor<0xi64>}> : (tensor<i32>) -> tensor<4xi32>