// 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.

  %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)

}
func rewriteValuePPC64_OpPPC64LessEqual(v *Value) bool {
	v_0 := v.Args[0]
	// match: (LessEqual (FlagEQ))
	// result: (MOVDconst [1])
	for {
		if v_0.Op != OpPPC64FlagEQ {
			break
		}
		v.reset(OpPPC64MOVDconst)
		v.AuxInt = int64ToAuxInt(1)
		return true
	}
	// match: (LessEqual (FlagLT))
	// result: (MOVDconst [1])
	for {
		if v_0.Op != OpPPC64FlagLT {
			break
		}

  %44 = "tf.Sub"(%42, %43) {device = ""} : (tensor<?xi64>, tensor<?xi64>) -> tensor<?xi64>
  %45 = "tf.LessEqual"(%10, %44) {device = ""} : (tensor<i64>, tensor<?xi64>) -> tensor<?xi1>
  %46 = "tf.All"(%45, %15) {device = "", keep_dims = false} : (tensor<?xi1>, tensor<1xi32>) -> tensor<i1>

	// that summary. But as soon as we need to iterate beyond that summary
	// in a level to find a large enough range, we'll stop narrowing.
	foundFree := func(addr offAddr, size uintptr) {
		if firstFree.base.lessEqual(addr) && addr.add(size-1).lessEqual(firstFree.bound) {
			// This range fits within the current firstFree window, so narrow
			// down the firstFree window to the base and bound of this range.
			firstFree.base = addr

		return true
	}
	return false
}
func rewriteValueARM64_OpARM64LessEqual(v *Value) bool {
	v_0 := v.Args[0]
	b := v.Block
	// match: (LessEqual (CMPconst [0] z:(AND x y)))
	// cond: z.Uses == 1
	// result: (LessEqual (TST x y))
	for {
		if v_0.Op != OpARM64CMPconst || auxIntToInt64(v_0.AuxInt) != 0 {
			break
		}
		z := v_0.Args[0]
		if z.Op != OpARM64AND {
			break