#         ; [condition-key] = map:{ [table] = list:{ [column] ; [column] } }
# }
#
# property-type: String, Number, Date, OrderBy, ...
# condition-key: NotEqual, GreaterThan, LessThan, GreaterEqual, LessEqual
#              , InScope, NotInScope, PrefixSearch, LikeSearch, NotLikeSearch
#              , EmptyString, FromTo, DateFromTo, RangeOf, ...
#              , (and prefix '!' means excluding, '%' means reviving)

func.func @preserve_assert(%arg0: tensor<1xi32>, %arg1: tensor<1xi32>) -> tensor<1xi1> {
  %0 = "tf.LessEqual"(%arg0, %arg1) : (tensor<1xi32>, tensor<1xi32>) -> tensor<1xi1>
  "tf.Assert"(%0, %arg1) {summarize = 3} : (tensor<1xi1>, tensor<1xi32>) -> ()
  func.return %0 : tensor<1xi1>
  // CHECK-LABEL: preserve_assert
  // CHECK: tfl.less_equal
  // CHECK: Assert
  // CHECK: return

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