// If b is strictly contained in a, thus forcing a split, it will throw.
func (a addrRange) subtract(b addrRange) addrRange {
	if b.base.lessEqual(a.base) && a.limit.lessEqual(b.limit) {
		return addrRange{}
	} else if a.base.lessThan(b.base) && b.limit.lessThan(a.limit) {
		throw("bad prune")
	} else if b.limit.lessThan(a.limit) && a.base.lessThan(b.limit) {
		a.base = b.limit

// Boolean-generating instructions (NOTE: NOT all boolean Values) always
// zero upper bit of the register; no need to zero-extend
(MOVBUreg x:((Equal|NotEqual|LessThan|LessThanU|LessThanF|LessEqual|LessEqualU|LessEqualF|GreaterThan|GreaterThanU|GreaterThanF|GreaterEqual|GreaterEqualU|GreaterEqualF) _)) => x

// omit unsigned extension
(MOVWUreg x) && zeroUpper32Bits(x, 3) => x

// don't extend after proper load

(Less64U x y) => (LessThanU (CMP x y))

(Leq8  x y) => (LessEqual (CMPW (SignExt8to32  x) (SignExt8to32  y)))
(Leq16 x y) => (LessEqual (CMPW (SignExt16to32 x) (SignExt16to32 y)))
(Leq32 x y) => (LessEqual (CMPW x y))
(Leq64 x y) => (LessEqual (CMP x y))

// Refer to the comments for op Less64F above.
(Leq32F x y) => (LessEqualF (FCMPS x y))
(Leq64F x y) => (LessEqualF (FCMPD x y))

(Less(32|64)U x y) => (LessThan ((CMPWU|CMPU) x y))

(Leq(8|16) x y)  => (LessEqual (CMPW (SignExt(8|16)to32 x) (SignExt(8|16)to32 y)))
(Leq(32|64) x y) => (LessEqual ((CMPW|CMP) x y))
(Leq(32|64)F x y) => (FLessEqual (FCMPU x y))

(Leq(8|16)U x y)  => (LessEqual (CMPWU (ZeroExt(8|16)to32 x) (ZeroExt(8|16)to32 y)))
(Leq(32|64)U x y) => (LessEqual (CMP(WU|U) x y))

// Absorb pseudo-ops into blocks.

 Const 4
 Floor 1
 Identity 1
 Mul 2
 Pow 1
 Sub 1
cluster 9 size 5
 All 1
 Const 1
 Less 1
 LogicalAnd 1
 LogicalNot 1
cluster 10 size 9
 All 1
 Const 4
 Equal 1
 LessEqual 1
 LogicalOr 1
 Max 1
cluster 11 size 302
 Add 24
 BatchMatMulV2 1
 BiasAdd 8
 Cast 8
 ConcatV2 16
 Const 81
 ExpandDims 3
 Fill 1
 GreaterEqual 8
 Identity 1
 Less 1

(Less32U x y) => (LessThanU (CMP x y))

(Leq8 x y)  => (LessEqual (CMP (SignExt8to32 x) (SignExt8to32 y)))
(Leq16 x y) => (LessEqual (CMP (SignExt16to32 x) (SignExt16to32 y)))
(Leq32 x y) => (LessEqual (CMP x y))
(Leq(32|64)F x y) => (GreaterEqual (CMP(F|D) y x)) // reverse operands to work around NaN

(Leq8U x y)  => (LessEqualU (CMP (ZeroExt8to32 x) (ZeroExt8to32 y)))
(Leq16U x y) => (LessEqualU (CMP (ZeroExt16to32 x) (ZeroExt16to32 y)))

  %0 = "tf.Less"(%arg0, %arg1) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi1>
  func.return %0: tensor<2xi1>
}

// CHECK-LABEL: func @less_equal
func.func @less_equal(%arg0: tensor<2xi32>, %arg1: tensor<2xi32>) -> tensor<2xi1> {
  // CHECK-NEXT:  mhlo.compare LE, %arg0, %arg1
  %0 = "tf.LessEqual"(%arg0, %arg1) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi1>
  func.return %0: tensor<2xi1>

		x := v_0
		if x.Op != OpARM64LessThanF {
			break
		}
		v.copyOf(x)
		return true
	}
	// match: (MOVBUreg x:(LessEqual _))
	// result: x
	for {
		x := v_0
		if x.Op != OpARM64LessEqual {
			break
		}
		v.copyOf(x)
		return true
	}
	// match: (MOVBUreg x:(LessEqualU _))
	// result: x
	for {
		x := v_0
		if x.Op != OpARM64LessEqualU {
			break
		}
		v.copyOf(x)

      %outputs_22, %control_23 = tf_executor.island wraps "tf.Const"() {device = "/device:CPU:0", value = dense<3> : tensor<i32>} : () -> tensor<i32>
      %outputs_24, %control_25 = tf_executor.island wraps "tf.LessEqual"(%outputs_0, %outputs_22) {device = "/device:CPU:0"} : (tensor<i32>, tensor<i32>) -> tensor<*xi1>

#include "tensorflow/cc/ops/standard_ops.h"

namespace tensorflow {
namespace ops {
namespace {

// Logical operations have no gradients.
REGISTER_NO_GRADIENT_OP("Less");
REGISTER_NO_GRADIENT_OP("LessEqual");
REGISTER_NO_GRADIENT_OP("Greater");
REGISTER_NO_GRADIENT_OP("GreaterEqual");
REGISTER_NO_GRADIENT_OP("Equal");
REGISTER_NO_GRADIENT_OP("ApproximateEqual");
REGISTER_NO_GRADIENT_OP("NotEqual");