diff --git a/crates/simd-masking-parity/src/lib.rs b/crates/simd-masking-parity/src/lib.rs index 52ba1bea..d231a0d1 100644 --- a/crates/simd-masking-parity/src/lib.rs +++ b/crates/simd-masking-parity/src/lib.rs @@ -35,7 +35,8 @@ //! the saturating narrows down the 512-bit ladder (`I32x16 → I16x32 → I8x64 / //! U8x64`) and the `F32x16` ↔ `I32x16` casts; `0x11xx` `F32x16` / `F64x8` //! `mul_add` rounding once, bit for bit against an exact integer reference (no -//! `mul_add`, so no libm in the binary). `main.rs` (native / qemu) and +//! `mul_add`, so no libm in the binary); `0x12xx` the `I32x16` gemmlowp requantization +//! methods against an independent `i128` reference. `main.rs` (native / qemu) and //! `selfcheck()` (the wasm cdylib export, driven by `run.mjs`) both call //! [`run`]. @@ -56,7 +57,7 @@ use ndarray::simd::{ }; /// Number of check groups [`run`] executes (for the log line only). -pub const CHECKS: usize = 18; +pub const CHECKS: usize = 19; /// The wasm export: identical to [`run`], `extern "C"` so `run.mjs` can call it. #[no_mangle] @@ -71,7 +72,7 @@ pub fn run() -> u32 { check_mask_algebra, check_care_match, check_masked_reductions, check_blend, check_morton_shift, check_predicates_under, check_set_range, check_unsigned_compare_to_mask, check_gather_scatter_group, check_f64x8_compare, check_i8x16_u8x16_lanes, check_care_match16, check_narrow_ladder, - check_mul_add_single_rounding, + check_mul_add_single_rounding, check_fixedpoint_requant, ]; for g in groups { if let Err(code) = g() { @@ -2077,6 +2078,161 @@ fn check_mul_add_single_rounding() -> Result<(), u32> { Ok(()) } +/// `I32x16::saturating_rounding_doubling_high_mul`, `rounding_divide_by_pot` and +/// `multiply_by_quantized_multiplier` against an independent `i128` reference +/// written from the definitions (no call into the scalar functions of +/// `simd_fixedpoint`): srdhm is `ab / 2^31` rounded half toward +inf with the +/// `MIN * MIN` saturation, rdpot is `x / 2^e` rounded half away from zero, and +/// the multiplier is their composition. Anti-vacuity: the rounding must differ +/// from truncation on some lanes, the saturating case must be hit, and exact ties +/// of both signs must occur. +fn check_fixedpoint_requant() -> Result<(), u32> { + fn o_srdhm(a: i32, b: i32) -> i32 { + let p = a as i128 * b as i128; + let q = (p + (1i128 << 30)).div_euclid(1i128 << 31); + q.clamp(i32::MIN as i128, i32::MAX as i128) as i32 + } + fn o_rdpot(x: i32, e: u32) -> i32 { + let d = 1i128 << e; + let m = (x as i128).abs(); + let r = (2 * m + d).div_euclid(2 * d); + (if x < 0 { -r } else { r }) as i32 + } + fn o_mbqm(x: i32, m: i32, s: i32) -> i32 { + let left = s.max(0) as u32; + let right = (-s).max(0) as u32; + let xs = ((x as i128) << left) as i32; // wraps like the realization's contract + o_rdpot(o_srdhm(xs, m), right) + } + + let mut rng = SplitMix64(0x1200_0000_0019); + const EDGE: [i32; 12] = + [i32::MIN, i32::MIN + 1, -1, 0, 1, i32::MAX - 1, i32::MAX, 1 << 30, -(1 << 30), (1 << 30) + 1, -(1 << 30) - 1, 1 << 15]; + const EXPS: [u32; 4] = [0, 1, 30, 31]; + let (mut trunc_differs_srdhm, mut trunc_differs_rdpot) = (0usize, 0usize); + let (mut min_min, mut tie_pos, mut tie_neg) = (0usize, 0usize, 0usize); + + for round in 0..512u32 { + // ---- M1: srdhm ---- + let mut a = [0i32; 16]; + let mut b = [0i32; 16]; + for i in 0..16 { + match (round as usize + i) % 4 { + 0 => { + a[i] = rng.next() as i32; + b[i] = rng.next() as i32; + } + 1 => { + a[i] = EDGE[(rng.next() % 12) as usize]; + b[i] = EDGE[(rng.next() % 12) as usize]; + } + 2 => { + // exact tie: odd * odd * 2^30 + let f = |r: &mut SplitMix64| { + let m = ((r.next() % 0x7FFF) as i32 | 1) << 15; + if r.next() & 1 == 0 { m } else { -m } + }; + a[i] = f(&mut rng); + b[i] = f(&mut rng); + } + _ => { + a[i] = if i % 2 == 0 { i32::MIN } else { rng.next() as i32 }; + b[i] = if i % 4 == 0 { i32::MIN } else { rng.next() as i32 }; + } + } + } + let got = I32x16::from_array(a).saturating_rounding_doubling_high_mul(I32x16::from_array(b)).to_array(); + for i in 0..16 { + let want = o_srdhm(a[i], b[i]); + if got[i] != want { + return Err(0x1200 + i as u32); + } + let p = a[i] as i128 * b[i] as i128; + if want as i128 != p / (1i128 << 31) { + trunc_differs_srdhm += 1; + } + if a[i] == i32::MIN && b[i] == i32::MIN { + min_min += 1; + } + if p.rem_euclid(1i128 << 31) == 1i128 << 30 { + if p >= 0 { tie_pos += 1 } else { tie_neg += 1 } + } + } + + // ---- M2: rounding_divide_by_pot ---- + let mut x = [0i32; 16]; + let mut e = [0i32; 16]; + for i in 0..16 { + let ex = match (round as usize + i) % 4 { + 0 => (rng.next() % 32) as u32, + 1 => EXPS[(rng.next() % 4) as usize], + 2 => EXPS[i % 4], + _ => 1 + (rng.next() % 31) as u32, + }; + let r = rng.next() as i32; + x[i] = match (round as usize + i) % 3 { + 0 => r, + 1 if ex >= 1 => ((r >> ex) << ex) | (1i32 << (ex - 1)), // exact tie + 1 => r, + _ => EDGE[(rng.next() % 12) as usize], + }; + e[i] = ex as i32; + } + let got = I32x16::from_array(x).rounding_divide_by_pot(I32x16::from_array(e)).to_array(); + for i in 0..16 { + let want = o_rdpot(x[i], e[i] as u32); + if got[i] != want { + return Err(0x1210 + i as u32); + } + if want as i128 != (x[i] as i128) / (1i128 << e[i]) { + trunc_differs_rdpot += 1; + } + } + + // ---- M3: multiply_by_quantized_multiplier ---- + let mut mx = [0i32; 16]; + let mut mm = [0i32; 16]; + let mut ms = [0i32; 16]; + for i in 0..16 { + let s = match (round as usize + i) % 4 { + 0 => (rng.next() % 62) as i32 - 31, + 1 => -31 + (rng.next() % 4) as i32, + 2 => 27 + (rng.next() % 4) as i32, + _ => (rng.next() % 5) as i32 - 2, + }; + ms[i] = s; + mm[i] = match (round as usize + i) % 3 { + 0 => (rng.next() >> 33) as i32 | (1 << 30), // normalised, as quantize_multiplier yields + 1 => (rng.next() >> 33) as i32, + _ => [0, 1, 1 << 30, i32::MAX][(rng.next() % 4) as usize], + }; + // keep x << max(s, 0) inside i32 (the documented contract) + mx[i] = (rng.next() as i32) >> (s.max(0) + 1); + } + let got = I32x16::from_array(mx) + .multiply_by_quantized_multiplier(I32x16::from_array(mm), I32x16::from_array(ms)) + .to_array(); + for i in 0..16 { + if got[i] != o_mbqm(mx[i], mm[i], ms[i]) { + return Err(0x1220 + i as u32); + } + } + } + if trunc_differs_srdhm < 256 { + return Err(0x12F0); + } + if trunc_differs_rdpot < 256 { + return Err(0x12F1); + } + if min_min == 0 { + return Err(0x12F2); + } + if tie_pos == 0 || tie_neg == 0 { + return Err(0x12F3); + } + Ok(()) +} + /// Exact `a * b + c`, rounded once to nearest-even: the reference for the /// `0x11xx` group. It holds the exact sum in a 4608-bit integer (bit 0 is /// `2^ORACLE_BASE`, below every product of two subnormals), so it shares no diff --git a/src/lib.rs b/src/lib.rs index d8cc364f..4bb6ae5c 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -490,6 +490,10 @@ pub mod simd_half; #[cfg(feature = "std")] pub(crate) mod simd_fma; +/// TFLite / gemmlowp fixed-point requantization (scalar reference + batch epilogue). +#[cfg(feature = "std")] +pub(crate) mod simd_fixedpoint; + /// Runtime SIMD dispatch — release-binary distribution path. /// /// Gated under `--features runtime-dispatch`. Mutually exclusive with diff --git a/src/simd.rs b/src/simd.rs index a84f5e6d..09de2045 100644 --- a/src/simd.rs +++ b/src/simd.rs @@ -997,6 +997,17 @@ pub use crate::simd_ops::{ scale_f32_inplace, sub_f32, sub_f32_inplace, }; +// Fixed-point requantization, bit-exact to TFLite-micro / gemmlowp (double rounding, +// TFLM's default build). The `I32x16` methods `saturating_rounding_doubling_high_mul`, +// `rounding_divide_by_pot` and `multiply_by_quantized_multiplier` are the vector form. +pub use crate::simd_fixedpoint::{ + multiply_by_quantized_multiplier_i32, multiply_by_quantized_multiplier_i64, + multiply_by_quantized_multiplier_single_rounding_i32, requantize_i32_to_i8, rounding_divide_by_pot_i32, srdhm_i16, + srdhm_i32, +}; +#[cfg(feature = "std")] +pub use crate::simd_fixedpoint::{quantize_multiplier, quantize_multiplier_single_rounding}; + // ChaCha20 keystream is NO LONGER an `ndarray::simd` surface: the AdaWorldAPI // `chacha20` fork (`vendor/chacha20/`) carries the accelerated backend, riding // the `U32x16` ARX lane above, and is `[patch]`ed transitively under the diff --git a/src/simd_avx.rs b/src/simd_avx.rs index 884c0a5c..084aef02 100644 --- a/src/simd_avx.rs +++ b/src/simd_avx.rs @@ -271,6 +271,32 @@ impl Shr for U64x8 { // --------------------------------------------------------------------------- impl I32x16 { + /// `SaturatingRoundingDoublingHighMul` per lane (TFLM `fixedpoint.h`). + /// AVX without AVX2 has no 256-bit integer multiply, so this arm goes + /// lane by lane through [`crate::simd_fixedpoint::srdhm_i32`]. + #[inline] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + let (a, b) = (self.to_array(), b.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::srdhm_i32(a[i], b[i]))) + } + + /// `RoundingDivideByPOT` per lane, per-lane exponent in `0..=31`. + #[inline] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + let (x, e) = (self.to_array(), exponent.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::rounding_divide_by_pot_i32(x[i], e[i]))) + } + + /// `MultiplyByQuantizedMultiplier` (double rounding) per lane, with + /// per-lane multiplier and shift. + #[inline] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + let (x, m, s) = (self.to_array(), multiplier.to_array(), shift.to_array()); + Self::from_array(core::array::from_fn(|i| { + crate::simd_fixedpoint::multiply_by_quantized_multiplier_i32(x[i], m[i], s[i]) + })) + } + /// Four 128-bit quarters of the 16 lanes. #[inline(always)] fn quarters(self) -> [__m128i; 4] { diff --git a/src/simd_avx2.rs b/src/simd_avx2.rs index ea1fb91d..20ed7997 100644 --- a/src/simd_avx2.rs +++ b/src/simd_avx2.rs @@ -2682,6 +2682,140 @@ impl U16x32 { } } +#[cfg(not(all(target_feature = "avx", not(target_feature = "avx2"))))] +impl I32x16 { + /// Store two 256-bit halves back into a fresh array (see `avx2_halves`). + #[inline(always)] + fn from_avx2_halves(lo: __m256i, hi: __m256i) -> Self { + let mut o = [0i32; 16]; + // SAFETY: two 32-byte `storeu` stores into the 64-byte local array; no alignment needed. + unsafe { + let p = o.as_mut_ptr() as *mut __m256i; + _mm256_storeu_si256(p, lo); + _mm256_storeu_si256(p.add(1), hi); + } + Self(o) + } + + /// `SaturatingRoundingDoublingHighMul` on one 256-bit half (exactness argument: see + /// [`Self::saturating_rounding_doubling_high_mul`]). + #[inline(always)] + fn srdhm_half(a: __m256i, b: __m256i) -> __m256i { + // SAFETY: AVX2 is a run-time precondition of this arm (see `avx2_halves`); register-only. + unsafe { + let nudge = _mm256_set1_epi64x(1 << 30); + let even = _mm256_srli_epi64(_mm256_add_epi64(_mm256_mul_epi32(a, b), nudge), 31); + let odd = _mm256_slli_epi64( + _mm256_add_epi64(_mm256_mul_epi32(_mm256_srli_epi64(a, 32), _mm256_srli_epi64(b, 32)), nudge), + 1, + ); + let r = _mm256_blend_epi32::<0xAA>(even, odd); + let min = _mm256_set1_epi32(i32::MIN); + let ovf = _mm256_and_si256(_mm256_cmpeq_epi32(a, min), _mm256_cmpeq_epi32(b, min)); + _mm256_blendv_epi8(r, _mm256_set1_epi32(i32::MAX), ovf) + } + } + + /// `RoundingDivideByPOT` on one 256-bit half, per-lane exponent. + #[inline(always)] + fn rdpot_half(x: __m256i, e: __m256i) -> __m256i { + // SAFETY: AVX2 run-time precondition (see `avx2_halves`); register-only. `vpsllvd` / + // `vpsravd` define all counts, so an out-of-contract exponent is unspecified, not UB. + unsafe { + let one = _mm256_set1_epi32(1); + let mask = _mm256_sub_epi32(_mm256_sllv_epi32(one, e), one); + let rem = _mm256_and_si256(x, mask); + let thr = _mm256_add_epi32(_mm256_srli_epi32(mask, 1), _mm256_srli_epi32(x, 31)); + let q = _mm256_srav_epi32(x, e); + // `cmpgt` lanes are -1 where rem > thr, so subtracting adds 1 there. + _mm256_sub_epi32(q, _mm256_cmpgt_epi32(rem, thr)) + } + } + + /// Saturating rounding doubling high multiply (gemmlowp / TFLM + /// `SaturatingRoundingDoublingHighMul`): per lane, `round(a * b / 2^31)` + /// with the half-way case rounded toward `+inf`, i.e. the C++ definition + /// `((a*b + nudge) / 2^31)` with `nudge = 1<<30` (product `>= 0`) or `1 - (1<<30)` + /// (product `< 0`) and **truncating** division. `i32::MIN * i32::MIN` saturates to + /// `i32::MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([3, -3, -1, 1, 0, i32::MIN, i32::MAX, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = a.saturating_rounding_doubling_high_mul(I32x16::splat(1 << 30)).to_array(); + /// // 1.5 -> 2, -1.5 -> -1 (tie toward +inf), -0.5 -> 0, 0.5 -> 1. + /// assert_eq!(r[..7], [2, -1, 0, 1, 0, -1_073_741_824, 1_073_741_824]); + /// let m = I32x16::splat(i32::MIN).saturating_rounding_doubling_high_mul(I32x16::splat(i32::MIN)); + /// assert!(m.to_array().iter().all(|&x| x == i32::MAX)); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + // Exact for both product signs: C++ truncating division after the sign-dependent nudge + // equals the flooring shift of `a*b + 2^30` (taking bits 31..=62 as the i32), computed + // with `vpmuldq` on even lanes and on odd lanes after a 32-bit shift down; `MIN*MIN` + // (2^31, not representable) is blended to `MAX`. + let (alo, ahi) = self.avx2_halves(); + let (blo, bhi) = b.avx2_halves(); + Self::from_avx2_halves(Self::srdhm_half(alo, blo), Self::srdhm_half(ahi, bhi)) + } + + /// Rounding arithmetic right shift (gemmlowp / TFLM `RoundingDivideByPOT`) with a + /// per-lane exponent: `x / 2^e` rounded to nearest, **ties away from zero**. + /// `exponent` lanes must be in `0..=31`; other values are a contract violation + /// (the result is unspecified, but no undefined behaviour). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([5, -5, 3, -3, 2, -2, 1, -1, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.rounding_divide_by_pot(I32x16::splat(1)).to_array(); + /// assert_eq!(r[..8], [3, -3, 2, -2, 1, -1, 1, -1]); // negative ties go away from zero + /// let e = I32x16::from_array([0, 1, 2, 3, 4, 5, 0, 0, 0, 0, 0, 0, 0, 0, 0, 31]); + /// let r = I32x16::splat(-10).rounding_divide_by_pot(e).to_array(); + /// assert_eq!(r[..6], [-10, -5, -3, -1, -1, 0]); + /// assert_eq!(r[15], 0); + /// let m = I32x16::splat(i32::MIN).rounding_divide_by_pot(I32x16::splat(31)).to_array(); + /// assert_eq!(m[0], -1); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + let (xlo, xhi) = self.avx2_halves(); + let (elo, ehi) = exponent.avx2_halves(); + Self::from_avx2_halves(Self::rdpot_half(xlo, elo), Self::rdpot_half(xhi, ehi)) + } + + /// Per-lane fixed-point requantization (gemmlowp / TFLM `MultiplyByQuantizedMultiplier`, + /// double rounding): with `left = max(shift, 0)`, `right = max(-shift, 0)`, + /// `rounding_divide_by_pot(srdhm(x << left, multiplier), right)`. + /// + /// Contract: `multiplier >= 0`, `-31 <= shift <= 30`, and `x << left` must not + /// overflow (TFLM asserts; this wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([100, -100, 7, -7, 7, -7, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let s = I32x16::from_array([0, 0, 0, 0, -1, -1, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.multiply_by_quantized_multiplier(I32x16::splat(1 << 30), s).to_array(); + /// assert_eq!(r[..7], [50, -50, 4, -3, 2, -2, 6]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + let (xlo, xhi) = self.avx2_halves(); + let (mlo, mhi) = multiplier.avx2_halves(); + let (slo, shi) = shift.avx2_halves(); + // SAFETY: AVX2 run-time precondition (see `avx2_halves`); register-only (`vpsllvd` wraps, + // like the scalar `wrapping_mul(1 << left)`). + unsafe { + let zero = _mm256_setzero_si256(); + let part = |x: __m256i, m: __m256i, s: __m256i| { + let left = _mm256_max_epi32(s, zero); + let right = _mm256_max_epi32(_mm256_sub_epi32(zero, s), zero); + Self::rdpot_half(Self::srdhm_half(_mm256_sllv_epi32(x, left), m), right) + }; + Self::from_avx2_halves(part(xlo, mlo, slo), part(xhi, mhi, shi)) + } + } +} + impl I32x16 { /// Narrow two 16×i32 vectors to one 32×i16 vector with signed saturation: lanes 0..16 /// come from `self`, lanes 16..32 from `hi`, each clamped to `i16::MIN..=i16::MAX`. diff --git a/src/simd_avx512.rs b/src/simd_avx512.rs index 794d9014..cc77d6e2 100644 --- a/src/simd_avx512.rs +++ b/src/simd_avx512.rs @@ -1030,6 +1030,105 @@ impl I32x16 { ) }) } + + /// Saturating rounding doubling high multiply (gemmlowp / TFLM + /// `SaturatingRoundingDoublingHighMul`): per lane, `round(a * b / 2^31)` + /// with the half-way case rounded toward `+inf`, i.e. the C++ definition + /// `((a*b + nudge) / 2^31)` with `nudge = 1<<30` (product `>= 0`) or `1 - (1<<30)` + /// (product `< 0`) and **truncating** division. `i32::MIN * i32::MIN` saturates to + /// `i32::MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([3, -3, -1, 1, 0, i32::MIN, i32::MAX, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = a.saturating_rounding_doubling_high_mul(I32x16::splat(1 << 30)).to_array(); + /// // 1.5 -> 2, -1.5 -> -1 (tie toward +inf), -0.5 -> 0, 0.5 -> 1. + /// assert_eq!(r[..7], [2, -1, 0, 1, 0, -1_073_741_824, 1_073_741_824]); + /// let m = I32x16::splat(i32::MIN).saturating_rounding_doubling_high_mul(I32x16::splat(i32::MIN)); + /// assert!(m.to_array().iter().all(|&x| x == i32::MAX)); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + // SAFETY: avx512f is compiled in on this arm; register-only ops. The exact result is + // `bits 31..=62 of (a*b + 2^30)` as an i32 for both product signs (C++ truncating + // division after the sign-dependent nudge equals a flooring shift by `+2^30`), taken from + // `vpmuldq` on the even lanes directly and on the odd lanes after a 32-bit shift down; + // only `MIN*MIN` (result 2^31) does not fit and is blended to `MAX`. + Self(unsafe { + let nudge = _mm512_set1_epi64(1 << 30); + let even = _mm512_srli_epi64(_mm512_add_epi64(_mm512_mul_epi32(self.0, b.0), nudge), 31); + let odd = _mm512_slli_epi64( + _mm512_add_epi64(_mm512_mul_epi32(_mm512_srli_epi64(self.0, 32), _mm512_srli_epi64(b.0, 32)), nudge), + 1, + ); + let r = _mm512_mask_blend_epi32(0xAAAA, even, odd); + let min = _mm512_set1_epi32(i32::MIN); + let ovf = _mm512_cmpeq_epi32_mask(self.0, min) & _mm512_cmpeq_epi32_mask(b.0, min); + _mm512_mask_blend_epi32(ovf, r, _mm512_set1_epi32(i32::MAX)) + }) + } + + /// Rounding arithmetic right shift (gemmlowp / TFLM `RoundingDivideByPOT`) with a + /// per-lane exponent: `x / 2^e` rounded to nearest, **ties away from zero**. + /// `exponent` lanes must be in `0..=31`; other values are a contract violation + /// (the result is unspecified, but no undefined behaviour). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([5, -5, 3, -3, 2, -2, 1, -1, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.rounding_divide_by_pot(I32x16::splat(1)).to_array(); + /// assert_eq!(r[..8], [3, -3, 2, -2, 1, -1, 1, -1]); // negative ties go away from zero + /// let e = I32x16::from_array([0, 1, 2, 3, 4, 5, 0, 0, 0, 0, 0, 0, 0, 0, 0, 31]); + /// let r = I32x16::splat(-10).rounding_divide_by_pot(e).to_array(); + /// assert_eq!(r[..6], [-10, -5, -3, -1, -1, 0]); + /// assert_eq!(r[15], 0); + /// let m = I32x16::splat(i32::MIN).rounding_divide_by_pot(I32x16::splat(31)).to_array(); + /// assert_eq!(m[0], -1); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + // SAFETY: avx512f is compiled in on this arm; register-only ops. `vpsllvd`/`vpsravd` + // define every shift count (counts above 31 give 0 / the sign), so a contract-violating + // exponent yields an unspecified value, never UB. + Self(unsafe { + let one = _mm512_set1_epi32(1); + let mask = _mm512_sub_epi32(_mm512_sllv_epi32(one, exponent.0), one); + let rem = _mm512_and_si512(self.0, mask); + let thr = _mm512_add_epi32(_mm512_srli_epi32(mask, 1), _mm512_srli_epi32(self.0, 31)); + let q = _mm512_srav_epi32(self.0, exponent.0); + let up = _mm512_cmpgt_epi32_mask(rem, thr); + _mm512_mask_add_epi32(q, up, q, one) + }) + } + + /// Per-lane fixed-point requantization (gemmlowp / TFLM `MultiplyByQuantizedMultiplier`, + /// double rounding): with `left = max(shift, 0)`, `right = max(-shift, 0)`, + /// `rounding_divide_by_pot(srdhm(x << left, multiplier), right)`. + /// + /// Contract: `multiplier >= 0`, `-31 <= shift <= 30`, and `x << left` must not + /// overflow (TFLM asserts; this wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([100, -100, 7, -7, 7, -7, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let s = I32x16::from_array([0, 0, 0, 0, -1, -1, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.multiply_by_quantized_multiplier(I32x16::splat(1 << 30), s).to_array(); + /// assert_eq!(r[..7], [50, -50, 4, -3, 2, -2, 6]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + // SAFETY: avx512f is compiled in on this arm; register-only ops (`vpsllvd` wraps, like + // the scalar `wrapping_mul(1 << left)`). + let (xs, right) = unsafe { + let zero = _mm512_setzero_si512(); + let left = _mm512_max_epi32(shift.0, zero); + let right = _mm512_max_epi32(_mm512_sub_epi32(zero, shift.0), zero); + (_mm512_sllv_epi32(self.0, left), right) + }; + Self(xs) + .saturating_rounding_doubling_high_mul(multiplier) + .rounding_divide_by_pot(Self(right)) + } } impl_bin_op!(I32x16, Add, add, _mm512_add_epi32); diff --git a/src/simd_fixedpoint.rs b/src/simd_fixedpoint.rs new file mode 100644 index 00000000..1a4de35f --- /dev/null +++ b/src/simd_fixedpoint.rs @@ -0,0 +1,460 @@ +//! gemmlowp / TFLite-Micro fixed-point requantization primitives. +//! +//! Scalar reference semantics (bit-exact with the TFLM *double-rounding* +//! default, i.e. `TFLITE_SINGLE_ROUNDING` undefined) for the integer +//! requantization epilogue of quantized FC/Conv kernels, plus a batch +//! epilogue ([`requantize_i32_to_i8`]) that rides the `I32x16` typed-wrapper +//! methods (`saturating_rounding_doubling_high_mul`, `rounding_divide_by_pot`, +//! `multiply_by_quantized_multiplier`) of whichever backend `crate::simd` +//! selected. +//! +//! Ported from (tflite-micro `tensorflow/lite/micro/kernels/internal/`): +//! +//! | here | TFLM | +//! |---|---| +//! | [`srdhm_i32`], [`srdhm_i16`] | `fixedpoint.h:319-348` `SaturatingRoundingDoublingHighMul` | +//! | [`rounding_divide_by_pot_i32`] | `fixedpoint.h:352-366` `RoundingDivideByPOT` | +//! | [`multiply_by_quantized_multiplier_i32`] | `common.cc:85-92` (double-rounding branch) | +//! | [`multiply_by_quantized_multiplier_single_rounding_i32`] | `common.cc:30-45` | +//! | [`multiply_by_quantized_multiplier_i64`] | `common.cc:47-110` (identical in both modes) | +//! | [`quantize_multiplier`] | `quantization_util.cc:54-106` `QuantizeMultiplier` | +//! | [`requantize_i32_to_i8`] | `reference/integer_ops/fully_connected.h:60-80,165-180` epilogue | + +/// `SaturatingRoundingDoublingHighMul` for `i32` (the NEON `VQRDMULH` semantics). +/// +/// Computes `round(2·a·b / 2^32)` with the C++ truncating division after a +/// sign-dependent nudge; `i32::MIN · i32::MIN` saturates to `i32::MAX`. +/// +/// ``` +/// use ndarray::simd::srdhm_i32; +/// assert_eq!(srdhm_i32(i32::MIN, i32::MIN), i32::MAX); +/// assert_eq!(srdhm_i32(1 << 15, 1 << 15), 1); // +0.5 rounds up +/// assert_eq!(srdhm_i32(-(1 << 15), 1 << 15), 0); // -0.5 rounds toward +inf +/// ``` +#[inline(always)] +pub fn srdhm_i32(a: i32, b: i32) -> i32 { + if a == b && a == i32::MIN { + return i32::MAX; + } + let p = a as i64 * b as i64; + let nudge: i64 = if p >= 0 { 1 << 30 } else { 1 - (1 << 30) }; + // `/` truncates toward zero (C++ semantics); it is NOT an arithmetic shift. + ((p + nudge) / (1i64 << 31)) as i32 +} + +/// `SaturatingRoundingDoublingHighMul` for `i16` (scalar reference only). +/// +/// ``` +/// use ndarray::simd::srdhm_i16; +/// assert_eq!(srdhm_i16(i16::MIN, i16::MIN), i16::MAX); +/// assert_eq!(srdhm_i16(1 << 7, 1 << 7), 1); +/// ``` +#[inline(always)] +pub fn srdhm_i16(a: i16, b: i16) -> i16 { + if a == b && a == i16::MIN { + return i16::MAX; + } + let p = a as i32 * b as i32; + let nudge: i32 = if p >= 0 { 1 << 14 } else { 1 - (1 << 14) }; + ((p + nudge) / (1i32 << 15)) as i16 +} + +/// Correctly-rounded (ties away from zero) division by `2^e`, `e` in `0..=31`. +/// +/// `e` outside `0..=31` is a contract violation (`debug_assert`); in release it +/// is clamped to 31 so the call is never UB. +/// +/// ``` +/// use ndarray::simd::rounding_divide_by_pot_i32; +/// assert_eq!(rounding_divide_by_pot_i32(1, 1), 1); // 0.5 -> 1 +/// assert_eq!(rounding_divide_by_pot_i32(-1, 1), -1); // -0.5 -> -1 +/// assert_eq!(rounding_divide_by_pot_i32(-3, 1), -2); // -1.5 -> -2 +/// ``` +#[inline(always)] +pub fn rounding_divide_by_pot_i32(x: i32, e: i32) -> i32 { + debug_assert!((0..=31).contains(&e), "rounding_divide_by_pot exponent {e} outside 0..=31"); + let e = e.clamp(0, 31); + let mask = ((1i64 << e) - 1) as i32; + let rem = x & mask; + let thr = (mask >> 1) + (x < 0) as i32; + (x >> e) + (rem > thr) as i32 +} + +/// TFLM `MultiplyByQuantizedMultiplier(int32, int32, int)`, double-rounding. +/// +/// Contract: `multiplier >= 0`, `-31 <= shift <= 30`, and `x << max(shift, 0)` +/// must not overflow. TFLM asserts the last point; this function **wraps**. +/// +/// ``` +/// use ndarray::simd::multiply_by_quantized_multiplier_i32; +/// // multiplier 0.5 in Q31 with shift 1 is a x1.0 scale. +/// assert_eq!(multiply_by_quantized_multiplier_i32(1000, 1 << 30, 1), 1000); +/// assert_eq!(multiply_by_quantized_multiplier_i32(1000, 1 << 30, -1), 250); +/// ``` +#[inline(always)] +pub fn multiply_by_quantized_multiplier_i32(x: i32, multiplier: i32, shift: i32) -> i32 { + debug_assert!(multiplier >= 0, "multiplier must be >= 0"); + debug_assert!((-31..=30).contains(&shift), "shift {shift} outside -31..=30"); + let left = shift.max(0) as u32; + let right = (-shift).max(0); + rounding_divide_by_pot_i32(srdhm_i32(x.wrapping_mul(1i32.wrapping_shl(left)), multiplier), right) +} + +/// TFLM single-rounding `MultiplyByQuantizedMultiplier(int32, int32, int)`. +/// +/// ``` +/// use ndarray::simd::multiply_by_quantized_multiplier_single_rounding_i32 as f; +/// assert_eq!(f(1000, 1 << 30, 1), 1000); +/// assert_eq!(f(1, 1 << 30, 0), 1); +/// assert_eq!(f(-1, 1 << 30, 0), 0); +/// ``` +#[inline(always)] +pub fn multiply_by_quantized_multiplier_single_rounding_i32(x: i32, multiplier: i32, shift: i32) -> i32 { + debug_assert!(multiplier >= 0, "multiplier must be >= 0"); + debug_assert!((-31..=30).contains(&shift), "shift {shift} outside -31..=30"); + let total = (31 - shift).clamp(1, 62); + let round = 1i64 << (total - 1); + ((x as i64 * multiplier as i64 + round) >> total) as i32 +} + +/// TFLM `MultiplyByQuantizedMultiplier(int64, int32, int)` (same in both +/// rounding modes). Used by the int16-activation / int64-accumulator kernels. +/// +/// Contract: `multiplier >= 0`, `-31 <= shift < 8`, `-(1<<47) <= x < (1<<47)`. +/// +/// ``` +/// use ndarray::simd::multiply_by_quantized_multiplier_i64; +/// assert_eq!(multiply_by_quantized_multiplier_i64(1 << 20, 1 << 30, 0), 1 << 19); +/// ``` +#[inline(always)] +pub fn multiply_by_quantized_multiplier_i64(x: i64, multiplier: i32, shift: i32) -> i32 { + debug_assert!(multiplier >= 0, "multiplier must be >= 0"); + debug_assert!((-31..8).contains(&shift), "shift {shift} outside -31..8"); + debug_assert!((-(1i64 << 47)..(1i64 << 47)).contains(&x), "x outside +-2^47"); + let reduced: i32 = if multiplier < 0x7FFF_0000 { + (multiplier + (1 << 15)) >> 16 + } else { + 0x7FFF + }; + let total = (15 - shift).clamp(1, 62); + ((x.wrapping_mul(reduced as i64) + .wrapping_add(1i64 << (total - 1))) + >> total) as i32 +} + +/// `frexp` for finite non-zero `f64`: `(q, e)` with `m = q * 2^e`, `|q|` in `[0.5, 1)`. +#[cfg(feature = "std")] +fn frexp(m: f64) -> (f64, i32) { + debug_assert!(m.is_finite() && m != 0.0); + let mut bits = m.to_bits(); + let mut adjust = 0i32; + if (bits >> 52) & 0x7FF == 0 { + // Subnormal: pre-scale by 2^54 so the exponent field is populated. + bits = (m * (1u64 << 54) as f64).to_bits(); + adjust = -54; + } + let exp_field = ((bits >> 52) & 0x7FF) as i32; + let q = f64::from_bits((bits & !(0x7FFu64 << 52)) | (1022u64 << 52)); + (q, exp_field - 1022 + adjust) +} + +#[cfg(feature = "std")] +fn quantize_multiplier_core(m: f64) -> (i64, i32) { + assert!(m.is_finite(), "quantize_multiplier: non-finite multiplier"); + if m == 0.0 { + return (0, 0); + } + let (q, mut shift) = frexp(m); + // `round` is round-half-away-from-zero, the same as TfLiteRound / std::round. + let mut q_fixed = (q * (1i64 << 31) as f64).round() as i64; + assert!(q_fixed <= (1i64 << 31), "quantize_multiplier: q_fixed overflow"); + if q_fixed == (1i64 << 31) { + q_fixed /= 2; + shift += 1; + } + assert!(q_fixed <= i32::MAX as i64); + // A shift below -31 would shift every bit out; flush to zero as TFLM does. + if shift < -31 { + shift = 0; + q_fixed = 0; + } + (q_fixed, shift) +} + +/// TFLM `QuantizeMultiplier` (double-rounding build): split a real multiplier +/// into a Q31 `(multiplier, shift)` pair so `m ~= multiplier * 2^(shift - 31)`. +/// +/// `0.0` maps to `(0, 0)`; magnitudes below about `2^-32` flush to `(0, 0)`. +/// +/// ``` +/// use ndarray::simd::quantize_multiplier; +/// assert_eq!(quantize_multiplier(0.5), (1 << 30, 0)); +/// assert_eq!(quantize_multiplier(1.0), (1 << 30, 1)); +/// assert_eq!(quantize_multiplier(0.0), (0, 0)); +/// ``` +#[cfg(feature = "std")] +pub fn quantize_multiplier(m: f64) -> (i32, i32) { + let (q, shift) = quantize_multiplier_core(m); + (q as i32, shift) +} + +/// TFLM `QuantizeMultiplier` for the `TFLITE_SINGLE_ROUNDING` build: as +/// [`quantize_multiplier`], and additionally saturates `shift > 30` to +/// `(i32::MAX, 30)`. +/// +/// ``` +/// use ndarray::simd::quantize_multiplier_single_rounding; +/// assert_eq!(quantize_multiplier_single_rounding(2f64.powi(40)), (i32::MAX, 30)); +/// ``` +#[cfg(feature = "std")] +pub fn quantize_multiplier_single_rounding(m: f64) -> (i32, i32) { + let (mut q, mut shift) = quantize_multiplier_core(m); + if shift > 30 { + shift = 30; + q = (1i64 << 31) - 1; + } + (q as i32, shift) +} + +/// Per-tensor / per-channel requantization epilogue `i32 -> i8`. +/// +/// `channels = multiplier.len() = shift.len() >= 1`; element `i` uses channel +/// `i % channels` (per-tensor = one channel). For each element: +/// `v = multiply_by_quantized_multiplier(acc, m, s) + output_offset` (wrapping), +/// clamped to `[act_min, act_max]`, stored as `i8`. +/// +/// Panics (never saturates silently) if the slice lengths are inconsistent, +/// `acc.len() % channels != 0`, or `act_min > act_max` / either is outside the +/// `i8` range. +/// +/// ``` +/// use ndarray::simd::requantize_i32_to_i8; +/// let acc = [100, -100, 0, 1_000_000]; +/// let mut out = [0i8; 4]; +/// requantize_i32_to_i8(&acc, &[1 << 30], &[0], 3, -128, 127, &mut out); +/// assert_eq!(out, [53, -47, 3, 127]); +/// ``` +pub fn requantize_i32_to_i8( + acc: &[i32], multiplier: &[i32], shift: &[i32], output_offset: i32, act_min: i32, act_max: i32, out: &mut [i8], +) { + use crate::simd::I32x16; + let channels = multiplier.len(); + assert!(channels >= 1 && shift.len() == channels, "multiplier/shift must have equal non-zero length"); + assert_eq!(acc.len(), out.len(), "acc/out length mismatch"); + assert!(acc.len().is_multiple_of(channels), "acc.len() must be a multiple of the channel count"); + assert!( + act_min <= act_max && act_min >= i8::MIN as i32 && act_max <= i8::MAX as i32, + "activation range must satisfy -128 <= act_min <= act_max <= 127" + ); + + // Channel pattern is identical for every 16-block when 16 % channels == 0. + let invariant = 16 % channels == 0; + let mut mv = [0i32; 16]; + let mut sv = [0i32; 16]; + for k in 0..16 { + mv[k] = multiplier[k % channels]; + sv[k] = shift[k % channels]; + } + let off = I32x16::splat(output_offset); + let lo = I32x16::splat(act_min); + let hi = I32x16::splat(act_max); + + let n16 = acc.len() / 16 * 16; + let mut base = 0; + while base < n16 { + if !invariant { + for k in 0..16 { + mv[k] = multiplier[(base + k) % channels]; + sv[k] = shift[(base + k) % channels]; + } + } + let x = I32x16::from_array(core::array::from_fn(|k| acc[base + k])); + let v = x.multiply_by_quantized_multiplier(I32x16::from_array(mv), I32x16::from_array(sv)) + off; + let r = v.simd_max(lo).simd_min(hi).to_array(); + for k in 0..16 { + out[base + k] = r[k] as i8; + } + base += 16; + } + for i in n16..acc.len() { + let c = i % channels; + let v = multiply_by_quantized_multiplier_i32(acc[i], multiplier[c], shift[c]).wrapping_add(output_offset); + out[i] = v.max(act_min).min(act_max) as i8; + } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn srdhm_min_min_saturates() { + assert_eq!(srdhm_i32(i32::MIN, i32::MIN), i32::MAX); + assert_eq!(srdhm_i16(i16::MIN, i16::MIN), i16::MAX); + // MIN * MAX is not the saturating case. + assert_eq!(srdhm_i32(i32::MIN, i32::MAX), -(i32::MAX)); + assert_eq!(srdhm_i32(i32::MAX, i32::MAX), 0x7FFF_FFFE); + } + + #[test] + fn srdhm_negative_ties_round_toward_positive_infinity() { + // p = +-2^30 (+-0.5), +-3*2^30 (+-1.5) + assert_eq!(srdhm_i32(1 << 15, 1 << 15), 1); + assert_eq!(srdhm_i32(-(1 << 15), 1 << 15), 0); + assert_eq!(srdhm_i32(3 << 15, 1 << 15), 2); + assert_eq!(srdhm_i32(-(3 << 15), 1 << 15), -1); + assert_eq!(srdhm_i16(1 << 7, 1 << 7), 1); + assert_eq!(srdhm_i16(-(1 << 7), 1 << 7), 0); + } + + #[test] + fn rdpot_ties_both_signs_at_e1() { + assert_eq!(rounding_divide_by_pot_i32(1, 1), 1); + assert_eq!(rounding_divide_by_pot_i32(-1, 1), -1); + assert_eq!(rounding_divide_by_pot_i32(3, 1), 2); + assert_eq!(rounding_divide_by_pot_i32(-3, 1), -2); + assert_eq!(rounding_divide_by_pot_i32(0, 1), 0); + } + + #[test] + fn rdpot_ties_both_signs_at_e31() { + assert_eq!(rounding_divide_by_pot_i32(1 << 30, 31), 1); + assert_eq!(rounding_divide_by_pot_i32(-(1 << 30), 31), -1); + assert_eq!(rounding_divide_by_pot_i32(i32::MAX, 31), 1); + assert_eq!(rounding_divide_by_pot_i32(i32::MIN, 31), -1); + assert_eq!(rounding_divide_by_pot_i32((1 << 30) - 1, 31), 0); + assert_eq!(rounding_divide_by_pot_i32(-(1 << 30) + 1, 31), 0); + } + + #[test] + fn rdpot_e0_is_identity() { + for x in [i32::MIN, -7, 0, 7, i32::MAX] { + assert_eq!(rounding_divide_by_pot_i32(x, 0), x); + } + } + + #[test] + fn mbqm_basics_and_rounding_modes() { + assert_eq!(multiply_by_quantized_multiplier_i32(1000, 1 << 30, 1), 1000); + assert_eq!(multiply_by_quantized_multiplier_i32(1000, 1 << 30, -1), 250); + assert_eq!(multiply_by_quantized_multiplier_i32(1, 1 << 30, 0), 1); + assert_eq!(multiply_by_quantized_multiplier_i32(-1, 1 << 30, 0), 0); + assert_eq!(multiply_by_quantized_multiplier_single_rounding_i32(1000, 1 << 30, 1), 1000); + assert_eq!(multiply_by_quantized_multiplier_single_rounding_i32(1000, 1 << 30, -1), 250); + assert_eq!(multiply_by_quantized_multiplier_single_rounding_i32(1, 1 << 30, 0), 1); + assert_eq!(multiply_by_quantized_multiplier_single_rounding_i32(-1, 1 << 30, 0), 0); + assert_eq!(multiply_by_quantized_multiplier_i64(1 << 20, 1 << 30, 0), 1 << 19); + } + + #[cfg(feature = "std")] + fn check_reconstruction(m: f64) { + let (q, s) = quantize_multiplier(m); + assert!(q >= 1 << 30, "q {q} not normalised for {m}"); + let back = q as f64 / (1i64 << 31) as f64 * 2f64.powi(s); + assert!((back - m).abs() <= m.abs() * 2f64.powi(-30), "m={m} q={q} s={s} back={back}"); + } + + #[cfg(feature = "std")] + #[test] + fn quantize_multiplier_known_values() { + assert_eq!(quantize_multiplier(0.0), (0, 0)); + assert_eq!(quantize_multiplier(1.0), (1 << 30, 1)); + assert_eq!(quantize_multiplier(0.5), (1 << 30, 0)); + assert_eq!(quantize_multiplier(2f64.powi(-32)), (1 << 30, -31)); + // Below the representable range: flushed to zero. + assert_eq!(quantize_multiplier(2f64.powi(-33)), (0, 0)); + // Subnormal input flushes the same way (and must not panic). + assert_eq!(quantize_multiplier(f64::MIN_POSITIVE / 4.0), (0, 0)); + // 1e-10 has frexp exponent -33 < -31, so TFLM flushes it to zero + // (the golden fixture records `Q 1e-10 0 0`). + assert_eq!(quantize_multiplier(1e-10), (0, 0)); + // Exact value from the TFLM golden fixture. + assert_eq!(quantize_multiplier(3.7), (1_986_422_374, 2)); + check_reconstruction(3.7); + check_reconstruction(0.999_999_999_9); + // Rounding up to 2^31 must renormalise: q_fixed = 2^31 -> 2^30, shift + 1. + let just_below_one = 1.0 - 2f64.powi(-53); + assert_eq!(quantize_multiplier(just_below_one), (1 << 30, 1)); + } + + #[cfg(feature = "std")] + #[test] + fn quantize_multiplier_single_rounding_saturates_high_shift() { + assert_eq!(quantize_multiplier_single_rounding(1.0), (1 << 30, 1)); + assert_eq!(quantize_multiplier_single_rounding(2f64.powi(40)), (i32::MAX, 30)); + assert_eq!(quantize_multiplier_single_rounding(0.0), (0, 0)); + // The double-rounding variant keeps the large shift. + assert_eq!(quantize_multiplier(2f64.powi(40)), (1 << 30, 41)); + } + + fn reference_epilogue(acc: &[i32], m: &[i32], s: &[i32], off: i32, lo: i32, hi: i32) -> Vec { + acc.iter() + .enumerate() + .map(|(i, &a)| { + let c = i % m.len(); + multiply_by_quantized_multiplier_i32(a, m[c], s[c]) + .wrapping_add(off) + .max(lo) + .min(hi) as i8 + }) + .collect() + } + + fn pseudo(n: usize, seed: u32) -> Vec { + let mut st = seed; + (0..n) + .map(|_| { + st = st.wrapping_mul(1664525).wrapping_add(1013904223); + (st as i32) >> 12 // up to +-2^19 so left shifts cannot overflow + }) + .collect() + } + + #[test] + fn requantize_per_tensor_with_tail() { + for len in [0usize, 1, 15, 16, 17, 37, 64, 100] { + let acc = pseudo(len, 7 + len as u32); + let mut out = vec![0i8; len]; + requantize_i32_to_i8(&acc, &[1_518_500_250], &[-3], 5, -100, 100, &mut out); + assert_eq!(out, reference_epilogue(&acc, &[1_518_500_250], &[-3], 5, -100, 100), "len {len}"); + } + // Hand-checked values. + let mut out = [0i8; 4]; + requantize_i32_to_i8(&[100, -100, 0, 1_000_000], &[1 << 30], &[0], 3, -128, 127, &mut out); + assert_eq!(out, [53, -47, 3, 127]); + } + + #[test] + fn requantize_per_channel_with_tail() { + // 3 channels (does not divide 16), 39 elements: two full blocks + 7-element tail. + let m3 = [1 << 30, 1_518_500_250, 2_000_000_000]; + let s3 = [0, -2, -5]; + let acc = pseudo(39, 11); + let mut out = vec![0i8; 39]; + requantize_i32_to_i8(&acc, &m3, &s3, -7, -128, 127, &mut out); + assert_eq!(out, reference_epilogue(&acc, &m3, &s3, -7, -128, 127)); + // 4 channels (divides 16), 36 elements: two blocks + 4 tail. + let m4 = [1 << 30, 1_100_000_000, 1_900_000_000, 1_300_000_000]; + let s4 = [1, -1, -4, 0]; + let acc = pseudo(36, 13); + let mut out = vec![0i8; 36]; + requantize_i32_to_i8(&acc, &m4, &s4, 2, -50, 60, &mut out); + assert_eq!(out, reference_epilogue(&acc, &m4, &s4, 2, -50, 60)); + // Clamping is exercised, not vacuous. + assert!(out.contains(&-50) || out.contains(&60)); + } + + #[test] + #[should_panic] + fn requantize_rejects_bad_activation_range() { + let mut out = [0i8; 1]; + requantize_i32_to_i8(&[1], &[1 << 30], &[0], 0, -200, 127, &mut out); + } + + #[test] + #[should_panic] + fn requantize_rejects_ragged_channels() { + let mut out = [0i8; 5]; + requantize_i32_to_i8(&[1; 5], &[1 << 30, 1 << 30], &[0, 0], 0, -128, 127, &mut out); + } +} diff --git a/src/simd_neon.rs b/src/simd_neon.rs index 699c9fbf..bfc8a81f 100644 --- a/src/simd_neon.rs +++ b/src/simd_neon.rs @@ -3267,10 +3267,101 @@ fn quad_mask4(cmp: uint32x4_t) -> u16 { unsafe { vaddvq_u32(vandq_u32(cmp, vld1q_u32(WEIGHTS.as_ptr()))) as u16 } } +/// gemmlowp NEON `RoundingDivideByPOT` on one quad, per-lane exponent: a rounding right shift +/// (`srshl` by `-e`, ties up) preceded by a `-1` fix-up on negative lanes (`x & -e` has its sign +/// bit set iff `x < 0 && e > 0`), which turns ties-up into ties-away-from-zero. Equals the scalar +/// `(x >> e) + (rem > thr)`: with `x = k*2^e + rem`, the fixed-up form is +/// `k + [rem - 1 + 2^(e-1) >= 2^e]` for `x < 0` and `k + [rem + 2^(e-1) >= 2^e]` otherwise. +#[cfg(target_arch = "aarch64")] +#[inline(always)] +fn rdpot_quad(x: int32x4_t, e: int32x4_t) -> int32x4_t { + // SAFETY: NEON baseline; pure register ops (`srshl` defines every shift count). + unsafe { + let neg = vnegq_s32(e); + let fixup = vshrq_n_s32::<31>(vandq_s32(x, neg)); + vrshlq_s32(vqaddq_s32(x, fixup), neg) + } +} + #[cfg(target_arch = "aarch64")] impl I32x16 { pub const LANES: usize = 16; + /// Saturating rounding doubling high multiply (gemmlowp / TFLM + /// `SaturatingRoundingDoublingHighMul`): per lane, `round(a * b / 2^31)` + /// with the half-way case rounded toward `+inf`, i.e. the C++ definition + /// `((a*b + nudge) / 2^31)` with `nudge = 1<<30` (product `>= 0`) or `1 - (1<<30)` + /// (product `< 0`) and **truncating** division. `i32::MIN * i32::MIN` saturates to + /// `i32::MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([3, -3, -1, 1, 0, i32::MIN, i32::MAX, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = a.saturating_rounding_doubling_high_mul(I32x16::splat(1 << 30)).to_array(); + /// // 1.5 -> 2, -1.5 -> -1 (tie toward +inf), -0.5 -> 0, 0.5 -> 1. + /// assert_eq!(r[..7], [2, -1, 0, 1, 0, -1_073_741_824, 1_073_741_824]); + /// let m = I32x16::splat(i32::MIN).saturating_rounding_doubling_high_mul(I32x16::splat(i32::MIN)); + /// assert!(m.to_array().iter().all(|&x| x == i32::MAX)); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + // SAFETY: NEON baseline; pure register ops. `sqrdmulh` is exactly this operation + // (round-half-up of `2ab / 2^32`, `MIN*MIN` saturating to `MAX`). + Self(core::array::from_fn(|p| I32x4(unsafe { vqrdmulhq_s32(self.0[p].0, b.0[p].0) }))) + } + + /// Rounding arithmetic right shift (gemmlowp / TFLM `RoundingDivideByPOT`) with a + /// per-lane exponent: `x / 2^e` rounded to nearest, **ties away from zero**. + /// `exponent` lanes must be in `0..=31`; other values are a contract violation + /// (the result is unspecified, but no undefined behaviour). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([5, -5, 3, -3, 2, -2, 1, -1, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.rounding_divide_by_pot(I32x16::splat(1)).to_array(); + /// assert_eq!(r[..8], [3, -3, 2, -2, 1, -1, 1, -1]); // negative ties go away from zero + /// let e = I32x16::from_array([0, 1, 2, 3, 4, 5, 0, 0, 0, 0, 0, 0, 0, 0, 0, 31]); + /// let r = I32x16::splat(-10).rounding_divide_by_pot(e).to_array(); + /// assert_eq!(r[..6], [-10, -5, -3, -1, -1, 0]); + /// assert_eq!(r[15], 0); + /// let m = I32x16::splat(i32::MIN).rounding_divide_by_pot(I32x16::splat(31)).to_array(); + /// assert_eq!(m[0], -1); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + Self(core::array::from_fn(|p| I32x4(rdpot_quad(self.0[p].0, exponent.0[p].0)))) + } + + /// Per-lane fixed-point requantization (gemmlowp / TFLM `MultiplyByQuantizedMultiplier`, + /// double rounding): with `left = max(shift, 0)`, `right = max(-shift, 0)`, + /// `rounding_divide_by_pot(srdhm(x << left, multiplier), right)`. + /// + /// Contract: `multiplier >= 0`, `-31 <= shift <= 30`, and `x << left` must not + /// overflow (TFLM asserts; this wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([100, -100, 7, -7, 7, -7, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let s = I32x16::from_array([0, 0, 0, 0, -1, -1, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let r = x.multiply_by_quantized_multiplier(I32x16::splat(1 << 30), s).to_array(); + /// assert_eq!(r[..7], [50, -50, 4, -3, 2, -2, 6]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + Self(core::array::from_fn(|p| { + // SAFETY: NEON baseline; pure register ops (`sshl` by a non-negative count wraps, + // like the scalar `wrapping_mul(1 << left)`). + I32x4(unsafe { + let zero = vdupq_n_s32(0); + let s = shift.0[p].0; + let left = vmaxq_s32(s, zero); + let right = vmaxq_s32(vnegq_s32(s), zero); + let xs = vshlq_s32(self.0[p].0, left); + rdpot_quad(vqrdmulhq_s32(xs, multiplier.0[p].0), right) + }) + })) + } + /// Narrow two 16×i32 vectors to one 32×i16 vector with signed saturation: lanes 0..16 /// come from `self`, lanes 16..32 from `hi`, each clamped to `i16::MIN..=i16::MAX`. /// diff --git a/src/simd_nightly/i_word_types.rs b/src/simd_nightly/i_word_types.rs index e70ad0bd..ed5f495c 100644 --- a/src/simd_nightly/i_word_types.rs +++ b/src/simd_nightly/i_word_types.rs @@ -355,6 +355,57 @@ pub struct I32x16(pub i32x16); impl I32x16 { pub const LANES: usize = 16; + /// Per-lane `SaturatingRoundingDoublingHighMul` (gemmlowp / NEON `VQRDMULH`): + /// `round(2*a*b / 2^32)`, ties toward +inf, `MIN * MIN` saturating to `MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([i32::MIN, 1 << 15, -(1 << 15), 3 << 15, 0, 1, -1, 7, 7, 7, 7, 7, 7, 7, 7, 7]); + /// let b = I32x16::from_array([i32::MIN, 1 << 15, 1 << 15, 1 << 15, 5, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1]); + /// let v = a.saturating_rounding_doubling_high_mul(b).to_array(); + /// assert_eq!(v[..7], [i32::MAX, 1, 0, 2, 0, 0, 0]); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + let (x, y) = (self.to_array(), b.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::srdhm_i32(x[i], y[i]))) + } + + /// Per-lane rounding arithmetic right shift (ties away from zero) by a + /// per-lane `exponent` in `0..=31`. Out-of-range exponents are a contract + /// violation (`debug_assert` in the scalar reference; never UB). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([1, -1, 3, -3, 1 << 30, -(1 << 30), 5, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let e = I32x16::from_array([1, 1, 1, 1, 31, 31, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.rounding_divide_by_pot(e).to_array()[..7], [1, -1, 2, -2, 1, -1, 5]); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + let (x, e) = (self.to_array(), exponent.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::rounding_divide_by_pot_i32(x[i], e[i]))) + } + + /// Per-lane TFLM `MultiplyByQuantizedMultiplier` (double rounding) with + /// per-lane `multiplier` (Q31, `>= 0`) and `shift` (`-31..=30`). `x << shift` + /// must not overflow (it wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::splat(1000); + /// let m = I32x16::splat(1 << 30); + /// let s = I32x16::from_array([1, -1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.multiply_by_quantized_multiplier(m, s).to_array()[..3], [1000, 250, 500]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + let (x, m, s) = (self.to_array(), multiplier.to_array(), shift.to_array()); + Self::from_array(core::array::from_fn(|i| { + crate::simd_fixedpoint::multiply_by_quantized_multiplier_i32(x[i], m[i], s[i]) + })) + } + /// Narrow two 16×i32 vectors to one 32×i16 vector with signed saturation: lanes 0..16 /// come from `self`, lanes 16..32 from `hi`, each clamped to `i16::MIN..=i16::MAX`. /// diff --git a/src/simd_scalar.rs b/src/simd_scalar.rs index a2461acb..6a049d3c 100644 --- a/src/simd_scalar.rs +++ b/src/simd_scalar.rs @@ -985,6 +985,57 @@ impl U16x32 { // Extra methods for I32x16 that float types have via the macro impl I32x16 { + /// Per-lane `SaturatingRoundingDoublingHighMul` (gemmlowp / NEON `VQRDMULH`): + /// `round(2*a*b / 2^32)`, ties toward +inf, `MIN * MIN` saturating to `MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([i32::MIN, 1 << 15, -(1 << 15), 3 << 15, 0, 1, -1, 7, 7, 7, 7, 7, 7, 7, 7, 7]); + /// let b = I32x16::from_array([i32::MIN, 1 << 15, 1 << 15, 1 << 15, 5, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1]); + /// let v = a.saturating_rounding_doubling_high_mul(b).to_array(); + /// assert_eq!(v[..7], [i32::MAX, 1, 0, 2, 0, 0, 0]); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + let (x, y) = (self.to_array(), b.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::srdhm_i32(x[i], y[i]))) + } + + /// Per-lane rounding arithmetic right shift (ties away from zero) by a + /// per-lane `exponent` in `0..=31`. Out-of-range exponents are a contract + /// violation (`debug_assert` in the scalar reference; never UB). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([1, -1, 3, -3, 1 << 30, -(1 << 30), 5, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let e = I32x16::from_array([1, 1, 1, 1, 31, 31, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.rounding_divide_by_pot(e).to_array()[..7], [1, -1, 2, -2, 1, -1, 5]); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + let (x, e) = (self.to_array(), exponent.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::rounding_divide_by_pot_i32(x[i], e[i]))) + } + + /// Per-lane TFLM `MultiplyByQuantizedMultiplier` (double rounding) with + /// per-lane `multiplier` (Q31, `>= 0`) and `shift` (`-31..=30`). `x << shift` + /// must not overflow (it wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::splat(1000); + /// let m = I32x16::splat(1 << 30); + /// let s = I32x16::from_array([1, -1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.multiply_by_quantized_multiplier(m, s).to_array()[..3], [1000, 250, 500]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + let (x, m, s) = (self.to_array(), multiplier.to_array(), shift.to_array()); + Self::from_array(core::array::from_fn(|i| { + crate::simd_fixedpoint::multiply_by_quantized_multiplier_i32(x[i], m[i], s[i]) + })) + } + /// Narrow two 16×i32 vectors to one 32×i16 vector with signed saturation: lanes 0..16 /// come from `self`, lanes 16..32 from `hi`, each clamped to `i16::MIN..=i16::MAX`. /// diff --git a/src/simd_wasm.rs b/src/simd_wasm.rs index 925cdfac..5b1274b7 100644 --- a/src/simd_wasm.rs +++ b/src/simd_wasm.rs @@ -2486,6 +2486,57 @@ pub mod wasm32_simd { impl I32x16 { pub const LANES: usize = 16; + /// Per-lane `SaturatingRoundingDoublingHighMul` (gemmlowp / NEON `VQRDMULH`): + /// `round(2*a*b / 2^32)`, ties toward +inf, `MIN * MIN` saturating to `MAX`. + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let a = I32x16::from_array([i32::MIN, 1 << 15, -(1 << 15), 3 << 15, 0, 1, -1, 7, 7, 7, 7, 7, 7, 7, 7, 7]); + /// let b = I32x16::from_array([i32::MIN, 1 << 15, 1 << 15, 1 << 15, 5, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1, 1]); + /// let v = a.saturating_rounding_doubling_high_mul(b).to_array(); + /// assert_eq!(v[..7], [i32::MAX, 1, 0, 2, 0, 0, 0]); + /// ``` + #[inline(always)] + pub fn saturating_rounding_doubling_high_mul(self, b: Self) -> Self { + let (x, y) = (self.to_array(), b.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::srdhm_i32(x[i], y[i]))) + } + + /// Per-lane rounding arithmetic right shift (ties away from zero) by a + /// per-lane `exponent` in `0..=31`. Out-of-range exponents are a contract + /// violation (`debug_assert` in the scalar reference; never UB). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::from_array([1, -1, 3, -3, 1 << 30, -(1 << 30), 5, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// let e = I32x16::from_array([1, 1, 1, 1, 31, 31, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.rounding_divide_by_pot(e).to_array()[..7], [1, -1, 2, -2, 1, -1, 5]); + /// ``` + #[inline(always)] + pub fn rounding_divide_by_pot(self, exponent: Self) -> Self { + let (x, e) = (self.to_array(), exponent.to_array()); + Self::from_array(core::array::from_fn(|i| crate::simd_fixedpoint::rounding_divide_by_pot_i32(x[i], e[i]))) + } + + /// Per-lane TFLM `MultiplyByQuantizedMultiplier` (double rounding) with + /// per-lane `multiplier` (Q31, `>= 0`) and `shift` (`-31..=30`). `x << shift` + /// must not overflow (it wraps). + /// + /// ``` + /// use ndarray::simd::I32x16; + /// let x = I32x16::splat(1000); + /// let m = I32x16::splat(1 << 30); + /// let s = I32x16::from_array([1, -1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]); + /// assert_eq!(x.multiply_by_quantized_multiplier(m, s).to_array()[..3], [1000, 250, 500]); + /// ``` + #[inline(always)] + pub fn multiply_by_quantized_multiplier(self, multiplier: Self, shift: Self) -> Self { + let (x, m, s) = (self.to_array(), multiplier.to_array(), shift.to_array()); + Self::from_array(core::array::from_fn(|i| { + crate::simd_fixedpoint::multiply_by_quantized_multiplier_i32(x[i], m[i], s[i]) + })) + } + /// Broadcast (`i32x4_splat` ×4). #[inline(always)] pub fn splat(v: i32) -> Self { diff --git a/tests/data/tflm_requant_golden.txt b/tests/data/tflm_requant_golden.txt new file mode 100644 index 00000000..0d61d239 --- /dev/null +++ b/tests/data/tflm_requant_golden.txt @@ -0,0 +1,2392 @@ +S -2147483648 -2147483648 2147483647 +S -2147483648 -2147483647 2147483647 +S -2147483648 -1073741824 1073741824 +S -2147483648 -2 2 +S -2147483648 -1 1 +S -2147483648 0 0 +S -2147483648 1 -1 +S -2147483648 2 -2 +S -2147483648 1073741824 -1073741824 +S -2147483648 2147483646 -2147483646 +S -2147483648 2147483647 -2147483647 +S -2147483647 -2147483648 2147483647 +S -2147483647 -2147483647 2147483646 +S -2147483647 -1073741824 1073741824 +S -2147483647 -2 2 +S -2147483647 -1 1 +S -2147483647 0 0 +S -2147483647 1 -1 +S -2147483647 2 -2 +S -2147483647 1073741824 -1073741823 +S -2147483647 2147483646 -2147483645 +S -2147483647 2147483647 -2147483646 +S -1073741824 -2147483648 1073741824 +S -1073741824 -2147483647 1073741824 +S -1073741824 -1073741824 536870912 +S -1073741824 -2 1 +S -1073741824 -1 1 +S -1073741824 0 0 +S -1073741824 1 0 +S -1073741824 2 -1 +S -1073741824 1073741824 -536870912 +S -1073741824 2147483646 -1073741823 +S -1073741824 2147483647 -1073741823 +S -2 -2147483648 2 +S -2 -2147483647 2 +S -2 -1073741824 1 +S -2 -2 0 +S -2 -1 0 +S -2 0 0 +S -2 1 0 +S -2 2 0 +S -2 1073741824 -1 +S -2 2147483646 -2 +S -2 2147483647 -2 +S -1 -2147483648 1 +S -1 -2147483647 1 +S -1 -1073741824 1 +S -1 -2 0 +S -1 -1 0 +S -1 0 0 +S -1 1 0 +S -1 2 0 +S -1 1073741824 0 +S -1 2147483646 -1 +S -1 2147483647 -1 +S 0 -2147483648 0 +S 0 -2147483647 0 +S 0 -1073741824 0 +S 0 -2 0 +S 0 -1 0 +S 0 0 0 +S 0 1 0 +S 0 2 0 +S 0 1073741824 0 +S 0 2147483646 0 +S 0 2147483647 0 +S 1 -2147483648 -1 +S 1 -2147483647 -1 +S 1 -1073741824 0 +S 1 -2 0 +S 1 -1 0 +S 1 0 0 +S 1 1 0 +S 1 2 0 +S 1 1073741824 1 +S 1 2147483646 1 +S 1 2147483647 1 +S 2 -2147483648 -2 +S 2 -2147483647 -2 +S 2 -1073741824 -1 +S 2 -2 0 +S 2 -1 0 +S 2 0 0 +S 2 1 0 +S 2 2 0 +S 2 1073741824 1 +S 2 2147483646 2 +S 2 2147483647 2 +S 1073741824 -2147483648 -1073741824 +S 1073741824 -2147483647 -1073741823 +S 1073741824 -1073741824 -536870912 +S 1073741824 -2 -1 +S 1073741824 -1 0 +S 1073741824 0 0 +S 1073741824 1 1 +S 1073741824 2 1 +S 1073741824 1073741824 536870912 +S 1073741824 2147483646 1073741823 +S 1073741824 2147483647 1073741824 +S 2147483646 -2147483648 -2147483646 +S 2147483646 -2147483647 -2147483645 +S 2147483646 -1073741824 -1073741823 +S 2147483646 -2 -2 +S 2147483646 -1 -1 +S 2147483646 0 0 +S 2147483646 1 1 +S 2147483646 2 2 +S 2147483646 1073741824 1073741823 +S 2147483646 2147483646 2147483644 +S 2147483646 2147483647 2147483645 +S 2147483647 -2147483648 -2147483647 +S 2147483647 -2147483647 -2147483646 +S 2147483647 -1073741824 -1073741823 +S 2147483647 -2 -2 +S 2147483647 -1 -1 +S 2147483647 0 0 +S 2147483647 1 1 +S 2147483647 2 2 +S 2147483647 1073741824 1073741824 +S 2147483647 2147483646 2147483645 +S 2147483647 2147483647 2147483646 +S 562021055 -551987228 -144461377 +S -2014677889 -298448725 279991910 +S -1923975998 1990036193 -1782915495 +S 1980641499 2105044278 1941499326 +S 1459427001 -749812279 -509571417 +S -602464986 -2125090481 596182703 +S 1550198204 -914196915 -659928851 +S -1801462009 2079828594 -1744708138 +S 932759010 -1052390830 -457105706 +S 1434261001 -512943701 -342584749 +S -1864201214 1646103662 -1428960098 +S 1739924932 -100604177 -81511082 +S 878259876 -1945440698 -795630043 +S -952245476 -1060101226 470074171 +S 1907130971 702647809 624005403 +S -1348764926 1944340367 -1221177211 +S 736059219 46394171 15901801 +S -1187669250 1516075123 -838467760 +S 702742283 1374253329 449710489 +S -1808571087 1508262631 -1270230946 +S -1338231143 962479708 -599781200 +S 459322474 776118705 166003017 +S 1731819792 1518530880 1224606220 +S 835991607 1956567717 761670149 +S -853242703 -2134146156 847943423 +S 1672538825 295671362 230279673 +S 1281484241 -1334121916 -796120712 +S -1403826006 -1613801610 1054954095 +S -220165227 1942138539 -199112749 +S 1825568144 -527639892 -448544779 +S 224678882 280114200 29306740 +S 232942829 -183281667 -19881013 +S -1379770758 -195812714 125810810 +S 1730701340 -1455371803 -1172914137 +S -344132971 1379087745 -220997987 +S 862280046 -804548194 -323050587 +S -1705047177 973586721 -773002994 +S 1500408664 -1892577266 -1322310105 +S -866150929 -696890749 281079006 +S -1526211907 -185142455 131580336 +S -1864686111 -661042258 573991014 +S 1838146061 -1695538763 -1451302272 +S 737868091 1013107420 348100270 +S 1669846801 -155411909 -120845660 +S -679053970 1551472478 -490589787 +S -1255009121 -1386489067 810276915 +S 998423488 -1311017253 -609527537 +S 435016043 1803343957 365303621 +S -1648334266 -350854404 269303721 +S -400036046 245531922 -45738006 +S -622617706 19604424 -5683890 +S -669318835 -1453929203 453154650 +S 150283905 -96223650 -6733865 +S -1352921390 -1790042698 1127732478 +S 562161143 1096989202 287166193 +S 1446765770 1550801323 1044779210 +S 640397408 -526926690 -157133903 +S -52533896 -1999480386 48913292 +S 451934762 537980376 113217176 +S 394764203 -877385683 -161286658 +S -154248641 316122116 -22706299 +S 1194523783 167067991 92930481 +S -245610924 -1148992389 131411982 +S 2042161202 537932408 511549736 +S -1672160090 248313305 -193351693 +S 417868262 1065946248 207417228 +S 687981449 -286320756 -91727529 +S 351212125 658719000 107730785 +S -1717084273 1923311578 -1537840843 +S 2028059303 -824591187 -778734604 +S 213411786 1085030363 107827721 +S -625579117 -678184934 197560681 +S -1597075689 1857469964 -1381393579 +S -1615564765 -1862624315 1401263389 +S -7269404 -2084799639 7057214 +S -365658942 -1651351733 281180966 +S -66090419 1853619680 -57046535 +S 290381831 125674060 16993593 +S -2000940326 877078559 -817227111 +S 1489300672 1605886362 1113697718 +S 1791401124 -611038293 -509719683 +S -662034412 610783089 -188294530 +S -210521473 -633550882 62108070 +S 1223174639 -2040547791 -1162265571 +S -1031332602 -595104143 285799757 +S -1907015722 -436439621 387568594 +S -1397092643 1228266529 -799075761 +S 1252920955 2107494868 1229590030 +S -887678468 -616163553 254695825 +S -1646790536 -6603827 5064122 +S -2018788253 -952447157 895368463 +S -1773260825 -192182091 158692232 +S -1642328839 -1077784083 824255814 +S 1822333940 -109695774 -93086777 +S -614560477 895001900 -256128979 +S 1683482263 2023908750 1586607882 +S -1151241911 -1488214786 797815264 +S 2013285810 -525205894 -492385390 +S 1353730933 1970316029 1242047994 +S 1031019720 787140749 377910973 +S -532642906 -927206115 229976028 +S -1281118923 972253734 -580014967 +S -644428763 -423692648 127144032 +S 1665128949 -1262455231 -978890225 +S -364119967 -1580608806 268002612 +S -455827876 1666071155 -353642588 +S -1940564234 245461868 -221810547 +S -1849035348 -413793885 356286541 +S -1341281801 1595255989 -996369787 +S 606418899 -501184589 -141527414 +S -1974422706 14951383 -13746484 +S -687774004 1191274462 -381529148 +S 335317043 416853412 65089228 +S 1902886908 1433751936 1270448690 +S -1405825213 -1950433593 1276828685 +S -393289055 407021444 -74541699 +S -1134487692 -413001766 218183464 +S -948970498 -534796200 236325812 +S 1407264039 -2025445952 -1327291714 +S 561993865 289518842 75766730 +S 2142939887 -540509502 -539365863 +S -482926673 -968297057 217750890 +S -297332995 -1928458935 267007607 +S -1964146065 488930963 -447189355 +S 167512035 -564868867 -44061958 +S -639162162 -1314671798 391289810 +S 1767803016 -3458676 -2847173 +S 1285613360 305082384 182640734 +S 1527565738 -935944554 -665763781 +S -1872820311 -250259985 218251712 +S -934267791 837910280 -364534830 +S -1342793799 -306391146 191582428 +S 878138426 982345098 401695714 +S -2103729699 667628073 -654025471 +S -807114609 -302331561 113628907 +S 488690030 1121148949 255133171 +S -61893009 346050504 -9973583 +S -942461230 -799565294 350903389 +S 318050332 -427916755 -63376066 +S 1560877208 -1069454805 -777322627 +S 1541935793 1243952021 893182189 +S 48015422 -9943191 -222319 +S -1974491457 -15091060 13875388 +S -1977459162 1649276698 -1518697160 +S 84695765 -1981566743 -78152079 +S -1304948798 -878639469 533917695 +S 427012630 1138859955 226454616 +S 1166862636 749210546 407093108 +S -129021357 -1313414357 78910264 +S -427943917 654873550 -130501181 +S 1423441648 1457577013 966142782 +S -725611810 10484017 -3542437 +S -812841191 2123954767 -803935305 +S 938319924 881471086 385149327 +S 1717124596 231050567 184747675 +S 1444727273 1087372607 731533794 +S -1618338458 -931119348 701689278 +S 1699437263 -629059211 -497813646 +S 905323740 -1147967030 -483953303 +S 995064535 -1019718099 -472499670 +S -1703056863 1301502406 -1032153426 +S -833226972 -777176411 301545648 +S 1274328841 -1291793728 -766557643 +S 727503648 361273368 122388682 +S 894451037 256620843 106885461 +S -1686358803 1869475242 -1468046583 +S 1174677059 -472773423 -258607834 +S -487655450 860121697 -195318383 +S 1427206414 551307621 366396165 +S 1787098291 1725668068 1436070751 +S 792873275 -23138889 -8543118 +S -1265333918 -2026993520 1194339084 +S 728634396 716837730 243220770 +S -1522986380 -1661453500 1178295841 +S 728355943 1958656723 664312050 +S 732140315 -824382512 -281056236 +S 2049615468 -1076946602 -1027866459 +S 568748104 747905635 198078301 +S 2131195267 -983573125 -976112852 +S 923003310 1665637795 715902633 +S -1976401094 367899582 -338590302 +S 1585219622 1730022138 1277059801 +S -528061669 -306612970 75395478 +S -1942168576 -1164639412 1053291405 +S 1875886390 1145183489 1000349466 +S -1894696841 1888846492 -1666504648 +S 1266794068 -462118298 -272602178 +S 60719300 652611715 18452353 +S -749884233 2130515782 -743959189 +S 800925979 -232597865 -86749751 +S -1817474704 -177152737 149929253 +S -446743384 1231028870 -256092289 +S -406241862 2016478975 -381459563 +S 1498208391 -517333085 -360921383 +S 536636160 1056348680 263971695 +S -1423469953 -142888064 94714046 +S 1944607464 -611551470 -553777233 +S -392694436 2128604399 -389242128 +S 216659144 -1664218625 -167902644 +S 1331518459 822456611 509953200 +S 1689513739 511855366 402697676 +S -1680004011 367972044 -287869251 +S 1697735112 -678101696 -536086531 +S 1405120254 -2113725098 -1383031694 +S 1481806737 -1101697155 -760193107 +S -907047465 -898754215 379613010 +S 318850567 811637099 120508927 +S 1320057177 -1003655297 -616946434 +S -2038993640 1727168550 -1639912691 +S 667495373 -994381519 -309080380 +S 1538148250 -1528370604 -1094704759 +S 621794969 -1941908292 -562271479 +S 970889200 -1865411389 -843362776 +S 1135029670 -290046776 -153301142 +S 184966377 1455965601 125404765 +S -380668274 -1221512407 216528317 +S -840858844 -1328116917 520031366 +S -1749830530 -19746524 16090027 +S -1491126754 -1185788171 823363879 +S -1426206978 -1480921328 983523359 +S 2075811648 -1688454578 -1632102616 +S -1782935243 -1202031301 997979180 +S 1291592142 729509383 438759377 +S 2055491822 -1230398849 -1177692214 +S -1467898423 -1155062854 789535670 +S -385218207 -210914722 37834137 +S 2082188344 -2060259403 -1997616195 +S 352716041 1431257108 235078549 +S -496345552 -1554882797 359378364 +S 1603446080 -1900722367 -1419198619 +S -1214736353 -1839477408 1040510870 +S -1987276382 -1745221956 1615024346 +S -1189127114 -301969010 167209440 +S -267409377 1419886076 -176807330 +S -29928968 1461910259 -20374295 +S -218750506 -1080684253 110082434 +S 437317143 447636781 91157499 +S 1944846881 1966496612 1780937799 +S -1395809113 944840714 -614122152 +S -226834733 -1468483466 155113197 +S -1782558596 1776601974 -1474701390 +S -1837903378 -188908598 161675620 +S 1839744670 -1986961464 -1702225657 +S 1192882820 1448213789 804452854 +S 927249678 -813623207 -351309710 +S -330832768 -1796369669 276741548 +S 1081460580 -1647657331 -829750883 +S -1772124902 774119105 -638810798 +S 814232777 -210603902 -79851877 +S -928948867 1060114614 -458579636 +S 2050690387 -35364119 -33770157 +S -1347117197 1437326484 -901635375 +S -1149303153 798962410 -427593485 +S -606101901 -1611569783 454846541 +S -181505419 479224767 -40504100 +S 305246010 -810843728 -115254341 +S 929475453 -1719796076 -744363403 +S 290329585 -978263109 -132256524 +S 515711232 12983801 3118018 +S 501539501 -1208773535 -282306073 +S 711179966 -2073083373 -686540903 +S -1913545781 -371872477 331362248 +S -2132248267 1974907491 -1960896456 +S -673384014 599124492 -187866788 +S -74882759 28328276 -987807 +S 1714951939 2013686929 1608103655 +S 852503474 1368004843 543067642 +S 727663632 1796748855 608818977 +S 1877462376 -176540987 -154342996 +S 1610209020 668728885 501420946 +S 1107915249 -1707878081 -881116916 +S -1712061662 -1947561467 1552675535 +S -1807757590 1689424882 -1422162472 +S 172410188 -79623157 -6392525 +S 952569245 805940055 357494554 +S 1109001757 50815456 26242076 +S -1844378449 1593083023 -1368228344 +S 607370329 -369503569 -104506269 +S -629430252 -1813828270 531635426 +S 161914098 1609683258 121365494 +S 1666800590 -216176570 -167788581 +S 20914335 -1773633541 -17273410 +S 700639711 552429433 180236063 +S -1404527842 -1421632707 929796471 +S 1130700018 -1268600343 -667947545 +S -58054523 -101432819 2742109 +S 1909409233 1837975398 1634213699 +S 1268616065 448047015 264681709 +S -1220516035 -2043298201 1161302542 +S -132147321 -65331780 4020249 +S -1644844730 -762393497 583948067 +S 1896141569 2043615632 1804430294 +S -518417143 -1819365602 439207217 +S 91602982 -767490076 -32738028 +S 1948359136 -758620262 -688277519 +S 164581200 -2139984778 -164006493 +S 2022457090 -1364904441 -1285439666 +S -1934896036 -432364979 389563518 +S 2061834644 4808462 4616684 +S 2036784966 312791643 296667831 +S -1925753754 -1597360911 1432431755 +S 2088608036 -1140119566 -1108861942 +S -591786288 -45110971 12431319 +S 451197751 1717010158 360752978 +S 1518374602 -484853054 -342814514 +S -1493307815 243805124 -169536144 +S -1067330973 926023529 -460247320 +S -323940729 978393191 -147587342 +S -422814448 -1516767859 298633876 +S -1708676970 1667955420 -1327133278 +S 554048914 87258649 22512656 +S -1472772043 866907945 -594536674 +S 942980662 -1937882175 -850942645 +S -1946889563 -893858703 810364345 +S -1626100408 -848104995 642195287 +S -1149012576 -1903389523 1018409850 +S 1726486693 1399510785 1125147914 +S 1160029273 275918321 149045759 +S -721571570 1308267274 -439588200 +S -123746783 -37436721 2157257 +S 2024471602 1136347646 1071255440 +S -1118932103 -544003075 283449192 +S 155174719 -66398124 -4797853 +S -1253653399 -1733958860 1012246785 +S 767301382 -1457083194 -520619540 +S -343776606 -1193449298 191051489 +S -280050080 1502604332 -195952348 +S 820019918 678870468 259227727 +S -1358126473 452727669 -286317166 +S -1136077334 148566123 -78595525 +S 1781998333 -1741789063 -1445349868 +S -1151653206 -1961872009 1052113338 +S -1605995718 -749766375 560712809 +S -1261459837 519343618 -305069199 +S -1833585498 1996350805 -1704543775 +S 982684003 2079406258 951531932 +S -104547760 565529083 -27532130 +S -662428243 1907473850 -588393095 +S 810418913 -1826033250 -689109732 +S 1818698991 -546944356 -463205924 +S -665915825 -1677414709 520151574 +S 689135050 -1336897176 -429015002 +S -1962177079 -2062508212 1884534181 +S 1972501775 -567413748 -521179580 +S 1499789287 2005934717 1400932390 +S -316764720 -833650057 122967608 +S -84381510 -2006392584 78837590 +S -1684686674 861690625 -675990578 +S -2121096652 209698880 -207122225 +S 152328570 1804129840 127973276 +S 972095478 -348058350 -157554610 +S -147064696 931864165 -63816235 +S 1337226529 1346395761 838393406 +S 125196404 1377935110 80332402 +S -1552983190 1062499030 -768361209 +S 176926675 853231566 70295960 +S -52215104 -89893189 2185713 +S 830073600 323525214 125053217 +S 914603443 1192814391 508014182 +S -566148065 942554479 -248488688 +S 223948217 158154593 16492996 +S 54926895 -836789957 -21402852 +S -2010671824 640875043 -600046195 +S -1608146356 312346266 -233900971 +S -82726413 1147389636 -44200303 +S 671645065 1783575172 557829376 +S -1300310141 -639815145 387410690 +S 1574985508 51540964 37800647 +S -1453288269 -1311873454 887797354 +S 846646337 -775331772 -305674879 +S 649813749 102617042 31051210 +S 1648632006 1831387026 1405963333 +S 2059052905 -421729523 -404363218 +S -213472199 320838658 -31893204 +S 251067044 632590607 73957561 +S -856609232 -2059727297 821604122 +S -591886900 2021272876 -557100836 +S -259334174 -1004517297 121307403 +S 1934949178 954366765 859913969 +S 1435651611 777001322 519446656 +S -234856253 502363080 -54940167 +S 1220266438 2097017806 1191590190 +S -917403912 -1184227943 505901569 +S 1566579830 1783411111 1300990523 +S -1524009828 -1119022140 794139104 +S -238840374 -515619386 57346526 +S -2123461912 1397336484 -1381705889 +S 711252971 -644804076 -213561028 +S -436738676 -888341281 180664004 +S -371434699 36321558 -6282277 +S 1551756744 1551010562 1120749442 +S 931933525 1415215015 614154300 +S -717605898 1829293879 -611279196 +S 781044144 -1804741391 -656388092 +S 2013762729 -881217652 -826345414 +S 530395824 -962551825 -237735672 +S 248993488 1995331242 231351930 +S 77900889 -1812285592 -65741436 +S -2136592343 -573136086 570229335 +S 1400576492 -237865432 -155134468 +H 12111 5397 1995 +H -25447 -14013 10882 +H -13965 -11032 4702 +H 29532 26600 23973 +H -24231 1008 -745 +H 1308 -7166 -286 +H -26598 -3880 3149 +H -14974 -18199 8316 +H -4233 4945 -639 +H 7450 -13728 -3121 +H -23655 8715 -6291 +H -19496 -11430 6801 +H -6602 31930 -6433 +H 2551 -20309 -1581 +H -7569 -8077 1866 +H -27407 -4099 3428 +H 14221 -14857 -6448 +H -29718 8687 -7878 +H -6551 -25000 4998 +H -14694 29596 -13272 +H -12178 25296 -9401 +H -3128 -3206 306 +H 3710 11614 1315 +H -28608 19618 -17127 +H -6520 27686 -5509 +H -29146 -25007 22243 +H -826 24744 -624 +H 26615 4880 3964 +H -17695 -30619 16535 +H 17264 30391 16012 +H -32579 16762 -16665 +H -5465 17998 -3002 +H -14278 -12425 5414 +H -29230 2854 -2546 +H -7540 -21023 4837 +H -15128 22493 -10384 +H 4042 -23335 -2878 +H -30515 -27553 25659 +H 32091 -6434 -6301 +H -1965 -8887 533 +H -14694 10537 -4725 +H -15628 -2427 1158 +H 27105 -9732 -8050 +H 3789 19654 2273 +H -6664 16599 -3376 +H 13574 10204 4227 +H 8771 22944 6141 +H -866 21291 -563 +H 10869 8300 2753 +H 1958 -32022 -1913 +H -16773 21827 -11173 +H 9336 -30870 -8795 +H -4671 -23363 3330 +H -20071 31947 -19568 +H -10127 -17589 5436 +H 12029 15188 5575 +H 15691 30886 14790 +H 17870 13089 7138 +H -372 -30802 350 +H -8074 -29106 7172 +H -6162 -1844 347 +H -19577 -15590 9314 +H -9491 26265 -7607 +H -9355 -10864 3102 +H -20754 -23144 14659 +H -7946 28312 -6865 +H 28001 -14007 -11969 +H -14710 21105 -9474 +H 3555 8592 932 +H -25490 18800 -14624 +H 22641 -8435 -5828 +H 20907 -6086 -3883 +H -29611 12387 -11194 +H 18920 25494 14720 +H -19025 19857 -11529 +H 23043 -11645 -8189 +H -13131 -5660 2268 +H -12390 -4214 1593 +H -26256 -11063 8864 +H 28896 -6204 -5471 +H -14857 6520 -2956 +H 23146 -20420 -14424 +H -182 20603 -114 +H -31791 20833 -20212 +H -23259 -20012 14205 +H 18561 27288 15457 +H 9629 4617 1357 +H -5514 -7305 1229 +H 15678 -1922 -920 +H 17636 609 328 +H 19774 28003 16899 +H -11240 2387 -819 +H 18278 -31184 -17394 +H -371 -24666 279 +H 18640 -19884 -11311 +H 32467 22446 22240 +H 15227 -12979 -6031 +H 9529 -13723 -3991 +H -30117 24675 -22679 +H 24604 -859 -645 +H 18984 -15922 -9224 +H -15450 -14214 6702 +H 28667 25383 22206 +H -24732 -4883 3685 +H -28026 -12059 10314 +H -7776 10648 -2527 +H 15931 25830 12558 +H 20140 -9656 -5935 +H 19111 -3926 -2290 +H 29817 -15064 -13707 +H 22771 7741 5379 +H -20018 -2613 1596 +H 32751 24782 24769 +H -11574 -22407 7914 +H -15424 4833 -2275 +H 13530 2424 1001 +H 3983 10424 1267 +H 20975 -14705 -9413 +H 14586 10571 4705 +H 5428 -25529 -4229 +H 25402 13010 10085 +H -11432 -20280 7075 +H 21786 -10320 -6861 +H -22691 -24650 17069 +H -27898 27303 -23245 +H -12043 -16554 6084 +H 7099 9879 2140 +H -20629 8392 -5283 +H -13586 -24931 10337 +H 27092 -32055 -26503 +H -23689 348 -252 +H -32425 6605 -6536 +H -14784 13059 -5892 +H 28546 16898 14721 +H 18739 -20262 -11587 +H 25948 28002 22174 +H -22625 5199 -3590 +H 4034 -29235 -3599 +H 16113 27228 13389 +H -1830 7605 -425 +H 32176 8811 8652 +H -67 -1391 3 +H -4612 -24178 3403 +H -88 -3866 10 +H 26387 12978 10451 +H 22075 -15893 -10707 +H -20996 -10629 6811 +H -471 6403 -92 +H 30163 27559 25368 +H -16499 21893 -11023 +H 8728 -7887 -2101 +H -8974 -712 195 +H -4753 -9467 1373 +H -7297 -13565 3021 +H -3015 -5474 504 +H 18458 19047 10729 +H 27774 -30420 -25784 +H -20886 -8263 5267 +H -12662 14168 -5475 +H 25518 -32507 -25315 +H 17072 -6100 -3178 +H 24771 -25362 -19172 +H -8836 27480 -7410 +H 13461 -15656 -6431 +H 31526 15315 14735 +H -23278 9463 -6722 +H 1868 15986 911 +H -14880 -1220 554 +H -22762 -26329 18289 +H -5310 -31847 5161 +H 30426 -31294 -29057 +H -5733 14408 -2521 +H 31100 27617 26211 +H 22268 10748 7304 +H 16308 -9877 -4916 +H 27994 -6247 -5337 +H 238 20427 148 +H -27498 -22570 18940 +H 9350 17908 5110 +H 3478 -16940 -1798 +H 21819 27501 18312 +H -15008 20704 -9483 +H 2717 -11153 -925 +H -5175 -9954 1572 +H -31470 -2822 2710 +H 17475 -11884 -6338 +H -9248 5873 -1658 +H 29837 19013 17312 +H -4274 12953 -1689 +H -31253 -32121 30636 +H 28758 -28879 -25345 +H 6151 14567 2734 +H -18418 -12741 7161 +H 3552 -25090 -2720 +H -11347 -10426 3610 +H 10348 22 7 +H 15244 -13130 -6108 +H -26086 9659 -7689 +H 13372 -21491 -8770 +H -15915 15422 -7490 +R -2147483648 0 -2147483648 +R -2147483647 0 -2147483647 +R -1073741824 0 -1073741824 +R -2 0 -2 +R -1 0 -1 +R 0 0 0 +R 1 0 1 +R 2 0 2 +R 1073741824 0 1073741824 +R 2147483646 0 2147483646 +R 2147483647 0 2147483647 +R -308617155 0 -308617155 +R 2068167175 0 2068167175 +R 1832532800 0 1832532800 +R 1704589082 0 1704589082 +R 424093582 0 424093582 +R 1436923664 0 1436923664 +R 252994604 0 252994604 +R -863019922 0 -863019922 +R -2147483648 1 -1073741824 +R -2147483647 1 -1073741824 +R -1073741824 1 -536870912 +R -2 1 -1 +R -1 1 -1 +R 0 1 0 +R 1 1 1 +R 2 1 1 +R 1073741824 1 536870912 +R 2147483646 1 1073741823 +R 2147483647 1 1073741824 +R -1707068342 1 -853534171 +R 681836417 1 340918209 +R 958009881 1 479004941 +R 1731876861 1 865938431 +R -1902975755 1 -951487878 +R -721117947 1 -360558974 +R 794700433 1 397350217 +R 463453171 1 231726586 +R -2147483648 2 -536870912 +R -2147483647 2 -536870912 +R -1073741824 2 -268435456 +R -2 2 -1 +R -1 2 0 +R 0 2 0 +R 1 2 0 +R 2 2 1 +R 1073741824 2 268435456 +R 2147483646 2 536870912 +R 2147483647 2 536870912 +R -840828072 2 -210207018 +R 299328070 2 74832018 +R 842403924 2 210600981 +R 109086998 2 27271750 +R 69261053 2 17315263 +R -1582018250 2 -395504563 +R 2049858992 2 512464748 +R 1569293918 2 392323480 +R -2147483648 3 -268435456 +R -2147483647 3 -268435456 +R -1073741824 3 -134217728 +R -2 3 0 +R -1 3 0 +R 0 3 0 +R 1 3 0 +R 2 3 0 +R 1073741824 3 134217728 +R 2147483646 3 268435456 +R 2147483647 3 268435456 +R 1555321543 3 194415193 +R -1491441636 3 -186430205 +R -739258638 3 -92407330 +R -396908324 3 -49613541 +R 280098149 3 35012269 +R 1201939564 3 150242446 +R 1691225034 3 211403129 +R -960603436 3 -120075430 +R -2147483648 4 -134217728 +R -2147483647 4 -134217728 +R -1073741824 4 -67108864 +R -2 4 0 +R -1 4 0 +R 0 4 0 +R 1 4 0 +R 2 4 0 +R 1073741824 4 67108864 +R 2147483646 4 134217728 +R 2147483647 4 134217728 +R 1515588266 4 94724267 +R 917581944 4 57348872 +R -972783823 4 -60798989 +R 661951480 4 41371968 +R -1814850811 4 -113428176 +R -2007743368 4 -125483961 +R -707277478 4 -44204842 +R 1422256760 4 88891048 +R -2147483648 5 -67108864 +R -2147483647 5 -67108864 +R -1073741824 5 -33554432 +R -2 5 0 +R -1 5 0 +R 0 5 0 +R 1 5 0 +R 2 5 0 +R 1073741824 5 33554432 +R 2147483646 5 67108864 +R 2147483647 5 67108864 +R -350499901 5 -10953122 +R 1573357296 5 49167416 +R -374216655 5 -11694270 +R -1262734992 5 -39460469 +R -564225996 5 -17632062 +R 877300976 5 27415656 +R 903219453 5 28225608 +R 648235056 5 20257346 +R -2147483648 6 -33554432 +R -2147483647 6 -33554432 +R -1073741824 6 -16777216 +R -2 6 0 +R -1 6 0 +R 0 6 0 +R 1 6 0 +R 2 6 0 +R 1073741824 6 16777216 +R 2147483646 6 33554432 +R 2147483647 6 33554432 +R -513815154 6 -8028362 +R 225209888 6 3518905 +R 729815047 6 11403360 +R 1511926624 6 23623854 +R 941834704 6 14716167 +R 1625479392 6 25398116 +R 540349868 6 8442967 +R 608601376 6 9509397 +R -2147483648 7 -16777216 +R -2147483647 7 -16777216 +R -1073741824 7 -8388608 +R -2 7 0 +R -1 7 0 +R 0 7 0 +R 1 7 0 +R 2 7 0 +R 1073741824 7 8388608 +R 2147483646 7 16777216 +R 2147483647 7 16777216 +R 1720141450 7 13438605 +R 269276352 7 2103722 +R -204469202 7 -1597416 +R 214453952 7 1675422 +R 1354665306 7 10583323 +R -2011852480 7 -15717598 +R -371935452 7 -2905746 +R -1767511488 7 -13808684 +R -2147483648 8 -8388608 +R -2147483647 8 -8388608 +R -1073741824 8 -4194304 +R -2 8 0 +R -1 8 0 +R 0 8 0 +R 1 8 0 +R 2 8 0 +R 1073741824 8 4194304 +R 2147483646 8 8388608 +R 2147483647 8 8388608 +R 715049497 8 2793162 +R 428195200 8 1672638 +R -1223489057 8 -4779254 +R 1571004288 8 6136736 +R -825709550 8 -3225428 +R -1110830720 8 -4339183 +R 1774142269 8 6930243 +R -1151936640 8 -4499753 +R -2147483648 9 -4194304 +R -2147483647 9 -4194304 +R -1073741824 9 -2097152 +R -2 9 0 +R -1 9 0 +R 0 9 0 +R 1 9 0 +R 2 9 0 +R 1073741824 9 2097152 +R 2147483646 9 4194304 +R 2147483647 9 4194304 +R 1343339982 9 2623711 +R -954120448 9 -1863517 +R 1213823808 9 2370750 +R 975672064 9 1905610 +R -1277635899 9 -2495383 +R 1893457664 9 3698160 +R -774356994 9 -1512416 +R -1901932288 9 -3714712 +R -2147483648 10 -2097152 +R -2147483647 10 -2097152 +R -1073741824 10 -1048576 +R -2 10 0 +R -1 10 0 +R 0 10 0 +R 1 10 0 +R 2 10 0 +R 1073741824 10 1048576 +R 2147483646 10 2097152 +R 2147483647 10 2097152 +R 377213571 10 368373 +R 560872960 10 547728 +R 1835408406 10 1792391 +R 1833326080 10 1790358 +R 1113995950 10 1087887 +R -1518673408 10 -1483080 +R -1518767667 10 -1483172 +R 1338213888 10 1306850 +R -2147483648 11 -1048576 +R -2147483647 11 -1048576 +R -1073741824 11 -524288 +R -2 11 0 +R -1 11 0 +R 0 11 0 +R 1 11 0 +R 2 11 0 +R 1073741824 11 524288 +R 2147483646 11 1048576 +R 2147483647 11 1048576 +R -45239736 11 -22090 +R 2072564736 11 1011995 +R 1489124303 11 727111 +R -613037056 11 -299335 +R 875917020 11 427694 +R 845675520 11 412928 +R -20900503 11 -10205 +R -995593216 11 -486130 +R -2147483648 12 -524288 +R -2147483647 12 -524288 +R -1073741824 12 -262144 +R -2 12 0 +R -1 12 0 +R 0 12 0 +R 1 12 0 +R 2 12 0 +R 1073741824 12 262144 +R 2147483646 12 524288 +R 2147483647 12 524288 +R -1877198413 12 -458300 +R 160733184 12 39242 +R -2110705289 12 -515309 +R -1010874368 12 -246796 +R -1055968517 12 -257805 +R -478160896 12 -116739 +R -876943993 12 -214098 +R -1791428608 12 -437361 +R -2147483648 13 -262144 +R -2147483647 13 -262144 +R -1073741824 13 -131072 +R -2 13 0 +R -1 13 0 +R 0 13 0 +R 1 13 0 +R 2 13 0 +R 1073741824 13 131072 +R 2147483646 13 262144 +R 2147483647 13 262144 +R 722019817 13 88137 +R 1951813632 13 238259 +R -309104089 13 -37732 +R 2047873024 13 249985 +R 665379832 13 81223 +R -527454208 13 -64387 +R -1543090048 13 -188365 +R -1200558080 13 -146553 +R -2147483648 14 -131072 +R -2147483647 14 -131072 +R -1073741824 14 -65536 +R -2 14 0 +R -1 14 0 +R 0 14 0 +R 1 14 0 +R 2 14 0 +R 1073741824 14 65536 +R 2147483646 14 131072 +R 2147483647 14 131072 +R -632179568 14 -38585 +R 1074978816 14 65612 +R 199606987 14 12183 +R 406642688 14 24820 +R 92762279 14 5662 +R -441655296 14 -26957 +R -1592944731 14 -97226 +R 408215552 14 24916 +R -2147483648 15 -65536 +R -2147483647 15 -65536 +R -1073741824 15 -32768 +R -2 15 0 +R -1 15 0 +R 0 15 0 +R 1 15 0 +R 2 15 0 +R 1073741824 15 32768 +R 2147483646 15 65536 +R 2147483647 15 65536 +R -1709604656 15 -52173 +R 987512832 15 30137 +R 1506297591 15 45969 +R 133513216 15 4075 +R -1087305975 15 -33182 +R 977846272 15 29842 +R 309749326 15 9453 +R -2115158016 15 -64550 +R -2147483648 16 -32768 +R -2147483647 16 -32768 +R -1073741824 16 -16384 +R -2 16 0 +R -1 16 0 +R 0 16 0 +R 1 16 0 +R 2 16 0 +R 1073741824 16 16384 +R 2147483646 16 32768 +R 2147483647 16 32768 +R 359807206 16 5490 +R -557809664 16 -8512 +R -1882749418 16 -28728 +R -1363312640 16 -20803 +R -1158951114 16 -17684 +R -889946112 16 -13580 +R 878600655 16 13406 +R -451313664 16 -6887 +R -2147483648 17 -16384 +R -2147483647 17 -16384 +R -1073741824 17 -8192 +R -2 17 0 +R -1 17 0 +R 0 17 0 +R 1 17 0 +R 2 17 0 +R 1073741824 17 8192 +R 2147483646 17 16384 +R 2147483647 17 16384 +R 1908594865 17 14561 +R -295632896 17 -2256 +R 851081447 17 6493 +R -962920448 17 -7347 +R -1871768645 17 -14280 +R 1263075328 17 9637 +R -2095625042 17 -15988 +R 1932591104 17 14745 +R -2147483648 18 -8192 +R -2147483647 18 -8192 +R -1073741824 18 -4096 +R -2 18 0 +R -1 18 0 +R 0 18 0 +R 1 18 0 +R 2 18 0 +R 1073741824 18 4096 +R 2147483646 18 8192 +R 2147483647 18 8192 +R -980807103 18 -3741 +R -1559363584 18 -5949 +R -1020447325 18 -3893 +R -1503789056 18 -5737 +R 701965598 18 2678 +R -337772544 18 -1289 +R -1331675641 18 -5080 +R 1871314944 18 7139 +R -2147483648 19 -4096 +R -2147483647 19 -4096 +R -1073741824 19 -2048 +R -2 19 0 +R -1 19 0 +R 0 19 0 +R 1 19 0 +R 2 19 0 +R 1073741824 19 2048 +R 2147483646 19 4096 +R 2147483647 19 4096 +R -882341750 19 -1683 +R 2000945152 19 3817 +R 1526312436 19 2911 +R -1746141184 19 -3331 +R -1642216371 19 -3132 +R -1170472960 19 -2233 +R -314243717 19 -599 +R -119275520 19 -228 +R -2147483648 20 -2048 +R -2147483647 20 -2048 +R -1073741824 20 -1024 +R -2 20 0 +R -1 20 0 +R 0 20 0 +R 1 20 0 +R 2 20 0 +R 1073741824 20 1024 +R 2147483646 20 2048 +R 2147483647 20 2048 +R -1634938852 20 -1559 +R -1980235776 20 -1889 +R -2113090460 20 -2015 +R 417857536 20 399 +R -1256375829 20 -1198 +R 1253572608 20 1196 +R 857217211 20 818 +R -2117599232 20 -2020 +R -2147483648 21 -1024 +R -2147483647 21 -1024 +R -1073741824 21 -512 +R -2 21 0 +R -1 21 0 +R 0 21 0 +R 1 21 0 +R 2 21 0 +R 1073741824 21 512 +R 2147483646 21 1024 +R 2147483647 21 1024 +R 325766837 21 155 +R 409993216 21 196 +R -1111424056 21 -530 +R -477102080 21 -228 +R -283107226 21 -135 +R -772800512 21 -369 +R -90085514 21 -43 +R -902823936 21 -431 +R -2147483648 22 -512 +R -2147483647 22 -512 +R -1073741824 22 -256 +R -2 22 0 +R -1 22 0 +R 0 22 0 +R 1 22 0 +R 2 22 0 +R 1073741824 22 256 +R 2147483646 22 512 +R 2147483647 22 512 +R 1539980680 22 367 +R 186646528 22 45 +R -823359725 22 -196 +R 698351616 22 167 +R 2130368728 22 508 +R -2116026368 22 -505 +R 1129100579 22 269 +R -1201668096 22 -287 +R -2147483648 23 -256 +R -2147483647 23 -256 +R -1073741824 23 -128 +R -2 23 0 +R -1 23 0 +R 0 23 0 +R 1 23 0 +R 2 23 0 +R 1073741824 23 128 +R 2147483646 23 256 +R 2147483647 23 256 +R 1224020247 23 146 +R 20971520 23 3 +R -554301741 23 -66 +R 2126512128 23 254 +R -2020890402 23 -241 +R 935329792 23 112 +R -1547204647 23 -184 +R 2042626048 23 244 +R -2147483648 24 -128 +R -2147483647 24 -128 +R -1073741824 24 -64 +R -2 24 0 +R -1 24 0 +R 0 24 0 +R 1 24 0 +R 2 24 0 +R 1073741824 24 64 +R 2147483646 24 128 +R 2147483647 24 128 +R 2133494483 24 127 +R 1333788672 24 80 +R 888510070 24 53 +R -1367343104 24 -82 +R 1084588931 24 65 +R -679477248 24 -41 +R 795768425 24 47 +R -92274688 24 -6 +R -2147483648 25 -64 +R -2147483647 25 -64 +R -1073741824 25 -32 +R -2 25 0 +R -1 25 0 +R 0 25 0 +R 1 25 0 +R 2 25 0 +R 1073741824 25 32 +R 2147483646 25 64 +R 2147483647 25 64 +R 2077762803 25 62 +R 1258291200 25 38 +R -2065625230 25 -62 +R 486539264 25 15 +R 1520728114 25 45 +R -754974720 25 -23 +R -136457571 25 -4 +R 1224736768 25 37 +R -2147483648 26 -32 +R -2147483647 26 -32 +R -1073741824 26 -16 +R -2 26 0 +R -1 26 0 +R 0 26 0 +R 1 26 0 +R 2 26 0 +R 1073741824 26 16 +R 2147483646 26 32 +R 2147483647 26 32 +R -275706856 26 -4 +R 503316480 26 8 +R -1132545404 26 -17 +R -1442840576 26 -22 +R -414224719 26 -6 +R 2113929216 26 32 +R 316293390 26 5 +R -1644167168 26 -25 +R -2147483648 27 -16 +R -2147483647 27 -16 +R -1073741824 27 -8 +R -2 27 0 +R -1 27 0 +R 0 27 0 +R 1 27 0 +R 2 27 0 +R 1073741824 27 8 +R 2147483646 27 16 +R 2147483647 27 16 +R 613971262 27 5 +R 201326592 27 2 +R -1497841054 27 -11 +R -1006632960 27 -8 +R -222006311 27 -2 +R -1811939328 27 -14 +R -613418385 27 -5 +R 1409286144 27 11 +R -2147483648 28 -8 +R -2147483647 28 -8 +R -1073741824 28 -4 +R -2 28 0 +R -1 28 0 +R 0 28 0 +R 1 28 0 +R 2 28 0 +R 1073741824 28 4 +R 2147483646 28 8 +R 2147483647 28 8 +R -918719402 28 -3 +R 939524096 28 4 +R 1594640424 28 6 +R -2013265920 28 -8 +R -1332710397 28 -5 +R -402653184 28 -2 +R -1249408360 28 -5 +R 2013265920 28 8 +R -2147483648 29 -4 +R -2147483647 29 -4 +R -1073741824 29 -2 +R -2 29 0 +R -1 29 0 +R 0 29 0 +R 1 29 0 +R 2 29 0 +R 1073741824 29 2 +R 2147483646 29 4 +R 2147483647 29 4 +R -1864290177 29 -3 +R 1879048192 29 4 +R -323439600 29 -1 +R 1879048192 29 4 +R 941433234 29 2 +R -805306368 29 -2 +R -395558678 29 -1 +R 268435456 29 1 +R -2147483648 30 -2 +R -2147483647 30 -2 +R -1073741824 30 -1 +R -2 30 0 +R -1 30 0 +R 0 30 0 +R 1 30 0 +R 2 30 0 +R 1073741824 30 1 +R 2147483646 30 2 +R 2147483647 30 2 +R 1617729709 30 2 +R 1610612736 30 2 +R -141967413 30 0 +R 536870912 30 1 +R 1729421682 30 2 +R 536870912 30 1 +R -1484728815 30 -1 +R 536870912 30 1 +R -2147483648 31 -1 +R -2147483647 31 -1 +R -1073741824 31 -1 +R -2 31 0 +R -1 31 0 +R 0 31 0 +R 1 31 0 +R 2 31 0 +R 1073741824 31 1 +R 2147483646 31 1 +R 2147483647 31 1 +R -1570920653 31 -1 +R -1073741824 31 -1 +R 1695968802 31 1 +R 1073741824 31 1 +R 1552853775 31 1 +R -1073741824 31 -1 +R -762268123 31 0 +R -1073741824 31 -1 +M -3 834802717 25 -39131377 +M 191522 526329557 11 96133890 +M 1221721685 1848374021 -7 8215279 +M -209565825 468348889 3 -365636950 +M 222 672313537 23 583021895 +M -336258357 1276972649 -3 -24993946 +M -1205 1602335569 20 -942780450 +M -12258460 1613066831 6 -589302636 +M -1097619426 2041887157 -29 -2 +M 1529406 1654641460 5 37709155 +M -7 802234591 28 -701955267 +M -10097 1048209730 16 -322991139 +M 1918118795 536320669 -30 0 +M -706334514 2016428262 -9 -1295369 +M 1419010 529652694 10 358382449 +M -1106204797 688145909 -9 -692335 +M 2082512102 588539511 -17 4354 +M -58291313 1766786490 -24 -3 +M -1711278714 1337546348 -1 -532929459 +M 1078755232 1728853159 -7 6784865 +M -591489316 1223419752 -23 -40 +M 1414458606 1793101891 -23 141 +M -1766526 669969597 3 -4408951 +M 923536050 2007096427 -6 13486904 +M 1 1221496226 27 76343514 +M 9027 1057700098 16 291377526 +M 894 65062776 16 1775089 +M -218788119 307176399 -29 0 +M 1 1411421594 28 176427699 +M -1739203315 362566483 -21 -140 +M -233 2095015758 23 -1906791686 +M -802400974 658251064 -24 -15 +M -77069489 1499959579 -18 -205 +M 36658 1896820187 15 1060999060 +M -231187571 1018644826 -6 -1713474 +M -16 957989362 24 -119748670 +M 1560239825 849330271 -8 2410450 +M -604093421 1680226664 -16 -7212 +M -1025870605 374573259 -13 -21843 +M -9 1165734539 27 -655725678 +M -1136134197 307700287 -28 -1 +M -50628965 2011629200 5 -1517634159 +M 2240169 431071044 9 230234144 +M -1530681286 429977528 -25 -9 +M 1569331646 1606989597 -16 17919 +M -23402 1545564354 12 -68987459 +M 38475 252117846 15 148013826 +M 1272185918 1896209155 0 1123328965 +M -938137495 637282867 -14 -16992 +M -916516111 1746319436 -6 -11645388 +M 8 2003734168 26 500933542 +M -13 1546737104 24 -157090487 +M -1392119380 1089567023 0 -706318472 +M -65 1135555818 22 -144162360 +M -115 1571620037 23 -706001188 +M 608 856403796 21 508489754 +M 196984053 2051420504 -1 94086194 +M -155577 249320553 13 -147966551 +M 1524 1481517670 13 8612949 +M -3 1412787328 23 -16556101 +M 32955899 2014098380 5 989085775 +M 56578643 1105336120 4 465947522 +M 207255750 1534961594 1 296281293 +M 1116449238 1596219125 -5 25932934 +M -1100755589 2073454919 -8 -4151602 +M 1312187742 1249497842 -25 23 +M -1317498792 1446791960 -19 -1693 +M 618487511 536006339 -20 147 +M -113239566 648663710 -14 -2088 +M 1689889438 1800172779 -22 338 +M 1286752090 230975539 -23 16 +M -2129692396 669369298 -17 -5065 +M -2086003307 139062263 -25 -4 +M 468305232 2016906464 -19 839 +M -5918339 1354810937 8 -955847550 +M 27482 392318757 16 329031497 +M -4993795 1888968682 -27 0 +M -1145021017 1489314207 -19 -1515 +M -5121 1486664644 16 -232336720 +M -1748029121 2005914988 -12 -398631 +M -915 17236373 20 -7700821 +M -686922723 1150806134 -29 -1 +M -29610 1090336765 12 -61578506 +M -1867 1719992256 18 -391995305 +M 2137480847 798677679 -19 1516 +M -88716440 981866010 -30 0 +M 1280809120 198067164 -20 113 +M 3 969364136 28 363511551 +M -583871813 940355681 0 -255670015 +M -1 306028389 30 -153014194 +M -1549956 845813988 10 -625121339 +M 58134438 1263116850 5 1094201031 +M -2098603080 2024262277 -22 -472 +M 9448697 1298307821 6 365594542 +M 447 1419359059 21 619583495 +M -32807 783815882 14 -196187192 +M -1704566439 56793738 -9 -88047 +M 2922650 1023955890 6 89188358 +M -423843414 1239303250 -3 -30574768 +M 390 1944542162 21 740597112 +M -2 1579581875 29 -789790937 +M 967239259 1096806722 -1 247004098 +M -13786870 2123550647 5 -436263035 +M -2134984814 1992921827 -20 -1890 +M -879752745 1700503354 0 -696639760 +M -740550762 1719147564 -26 -9 +M 1447708619 998049172 -6 10512918 +M 110949767 2026218706 2 418738450 +M -208400527 1153787923 -31 0 +M -836999406 1519153466 -17 -4517 +M -2298309 2096990162 8 -574532910 +M -465083783 880751829 1 -381491513 +M 1240141605 1961364592 -18 4321 +M 5422 1256001820 18 831303939 +M 1984665687 1311755164 -4 75768780 +M -300152136 977832919 2 -546683817 +M 1586482789 330199643 -14 14889 +M 1325184620 235861162 -22 35 +M -624417556 284624581 0 -82759459 +M -152 1268376981 23 -753098832 +M -119798562 1020869136 -12 -13904 +M 1294335293 401789598 0 242167365 +M -853 539788090 18 -56205962 +M 5849 1112398753 18 794240760 +M -689475092 467959111 -31 0 +M 1676872231 1696331812 -15 40423 +M -115965723 1546709475 3 -668187747 +M -335328683 1224782882 -13 -23346 +M -1823301412 1263740005 -7 -8382554 +M -31342355 1055653745 5 -493029869 +M 393679 352933448 12 265011762 +M -1640257078 1929822792 -23 -176 +M -278607374 558673395 -21 -35 +M 948574670 884080931 -16 5959 +M 259537 2025418488 13 2005275872 +M -674 819392818 21 -539326913 +M 1102598375 1907884290 -9 1913240 +M 0 670078803 30 0 +M 683594393 1475053213 -31 0 +M -20719 416975314 16 -263650865 +M 20514 527251518 15 165039637 +M 6 827668446 28 620751335 +M -3151 706718445 18 -271834695 +M -16879 32143948 16 -16557547 +M -545 1577639118 21 -839661445 +M -12765029 1687712742 7 -1284104711 +M -595642728 154938886 -26 -1 +M 698668479 1494052953 -20 464 +M 1363019923 1918583689 -20 1161 +M 0 903390293 30 0 +M -11780428 1328898259 6 -466555067 +M -16 4020184 27 -4020184 +M -1842 1285833132 20 -1156496401 +M -1698192290 655592540 -1 -259215524 +M 866581706 1419643798 1 1145747811 +M -360634741 1490223930 2 -1001034902 +M 773 1152844799 20 435131362 +M -61 1553175432 22 -185046292 +M -1681206533 2022556226 -21 -755 +M 381529235 747808229 -26 2 +M -1884188044 5095155 -18 -17 +M -1 700037833 27 -43752365 +M -204167882 1088646021 -25 -3 +M -14 1573142140 27 -1376499372 +M -2 2069210485 30 -2069210485 +M -45720 1274444390 13 -222273245 +M -931199657 476236659 -24 -12 +M 674014528 1303913112 -27 3 +M 625067872 2061659978 -18 2289 +M 957292605 1094955140 -16 7448 +M -754225202 1022868909 -16 -5482 +M -302 1586248973 22 -935639043 +M 4906 1510617354 18 904673918 +M -1783666474 697399678 -25 -17 +M 645920956 840006021 -4 15791083 +M 787936027 1271562329 -2 116637648 +M -3685216 1436323189 9 -1261987972 +M -1400764505 1003383742 -10 -639149 +M 4606 400208543 18 225019598 +M -144947 630861412 13 -348821522 +M -771503990 1051174419 -10 -368793 +M -2 483072434 30 -483072434 +M 1257231735 220086125 -19 246 +M -371609468 1845709130 -10 -311903 +M -11 885338448 27 -608670183 +M -109702 474447144 13 -198546603 +M 10657709 1689293214 7 1073121756 +M 1341118099 1098355785 -10 669854 +M 870290723 702772401 -15 8692 +M -983229070 689499867 -20 -301 +M -26477376 1102955991 6 -870328560 +M 95223027 279607308 3 99186056 +M -271711 1605631966 12 -832114920 +M 194609067 1018980887 -12 22544 +M -29247 1385996551 15 -618533953 +M -750505031 2135124616 -9 -1457394 +M 1476 1441994805 16 64953135 +M 9314 263841475 17 149988983 +M 387129552 1659193474 -21 143 +M 1528606455 61302462 -15 1332 +M -1012346087 201835864 -24 -6 +M 1811416785 587861973 -25 15 +M -1799178753 598017293 -6 -7828494 +M -15332 82005312 17 -76739834 +M 104817 1558785294 14 1246545396 +M 35 1957671191 17 4182037 +M -2 1872012987 29 -936006493 +M 2023792406 1590127878 -11 731709 +M 1745951263 356900880 -10 283367 +M 16748429 29776905 3 1857863 +M 262722414 1783843686 -30 0 +M 131675036 690234540 4 677158370 +M 798638493 1578971930 -19 1120 +M -258132425 784875226 -25 -3 +M 196776742 2129530235 3 1561053178 +M -1113709581 303937520 -2 -39406369 +M 913740761 2132629191 -24 54 +M -109856395 1744411101 -5 -2788652 +M -2051 1053482230 18 -263756354 +M 4899184 378299124 7 110468687 +M 256017 242356823 13 236692302 +M -629990293 1125283313 -16 -5037 +M -15995 1625393065 17 -1586801884 +M 2894 1197903365 17 211592550 +M -432350040 1602791290 -30 0 +M 1990 902235474 19 438341942 +M 1765552017 804837723 -21 316 +M 7 404041355 25 44192023 +M -1681587715 1499739332 -24 -70 +M -544628 173802973 11 -90272871 +M 254079249 1941039449 -27 2 +M -8 1893909545 28 -1893909545 +M -12 177334827 27 -133001120 +M -10605 1016562641 17 -657998462 +M 95266213 387847038 4 275289405 +M -31257482 1911038444 -8 -108656 +M 1762115662 2058292185 -30 2 +M 1415989642 335066851 -2 55233388 +M 254585659 1741298814 -23 25 +M -1090267751 596429753 -3 -37850587 +M -1685021145 297943242 -26 -3 +M -15 490263718 27 -459622236 +M 1684505322 840670683 -22 157 +M -2561 2009449124 19 -1256396291 +M 22033001 2123614238 -26 0 +M -641232529 1795793787 -19 -1023 +M 477552330 164990417 -17 280 +M -30823810 1290328700 6 -1185323199 +M 815896051 377481872 -2 35854286 +M 770 705672511 21 530632650 +M 2419767 590934087 6 42615020 +M 170734395 1167061407 1 185573030 +M -8168149 1704513651 4 -103732358 +M -682877002 738709318 -28 -1 +M 101671 1400753337 10 67909237 +M -9455 1468436860 17 -847416413 +M 118 1129630141 24 1041377786 +M 3472905 1528611458 9 1265698046 +M -1169455 1367171537 9 -381194494 +M 53850216 162844539 5 130671465 +M 1 122140849 30 61070425 +M 241699184 1221471367 -30 0 +M -1657213851 1706956234 -5 -41164334 +M -286964 1322303379 10 -180937513 +M 1294820619 655032423 -25 12 +M -99 1615613893 24 -1249576370 +M 2139816229 743869194 -24 44 +M -7013278 583414096 -14 -116 +M 31143721 892224150 6 828122497 +M 2860 774141377 19 540538169 +M 587436047 1650696862 -14 27560 +M 161102170 1393744532 3 836459058 +M -2 2138935339 29 -1069467669 +M 838627674 1144435253 -26 7 +M -717697701 1862832828 -22 -148 +M 117660682 1856235648 4 1627251150 +M 1565683 166742829 10 124486166 +M -1166446 1813530411 10 -1008694312 +M -29758697 522842993 -1 -3622642 +M 491824 1157614665 9 135741872 +M 778 399451511 21 303489527 +M -5128 1759627062 18 -1101485300 +M -1627844702 4818805 -11 -1784 +M -800405840 2036300548 -23 -90 +M -1212589490 1419058914 -23 -96 +M 38 760588991 25 451599713 +M -650117290 231902946 -30 0 +M -3486 1366521962 18 -581505806 +M 19 1480084947 22 54925027 +M -1346071397 1313525579 -11 -402019 +M 1324085333 1695215756 -24 62 +M 169194488 1364869430 -3 13441801 +M -308378808 925880920 -20 -127 +M 913800770 83020975 -18 135 +M -1848334756 361759985 -29 -1 +M 120292306 377412309 -19 40 +M 439849080 605717368 -5 3876983 +M -15 1755008714 27 -1645320669 +M 1990683474 1012963126 -30 1 +M -1 247666778 28 -30958347 +M 575483120 638605628 -25 5 +M 1 749887418 30 374943709 +M 15649 1134152036 16 541636512 +M 2067715993 179728153 -13 21125 +M 581547831 1728592647 -13 57142 +M 202797216 2109564236 -11 97274 +M 369901738 322116730 -24 3 +M 1965936923 1339823351 -11 598904 +M -473160814 1732507706 -2 -95432014 +M 684857929 140417498 -6 699700 +M 603691311 1495744771 -24 25 +M 1507763798 495581822 -22 83 +M 159875320 923130506 -3 8590623 +M -1289762236 1792326596 -22 -257 +M -554593158 1865466912 -29 -1 +M 901445975 438912749 -26 3 +M 107598176 604222788 -11 14782 +M 21 12892974 22 528813 +M 36 1577474299 24 443664647 +M -278377659 1081911783 -17 -1070 +M 1190460501 590089647 -25 10 +M -359297125 1906748030 -8 -1246170 +M -18155173 1278299764 6 -691644948 +M 1104930015 1623615839 -24 50 +M -279969652 1893484216 -16 -3767 +M 14 1622766907 26 709960522 +M 187 77307434 23 56470665 +M -1650 94326505 13 -593715 +M 1035004564 1452575434 -15 21365 +M -1982046026 994029781 -3 -114681477 +M 1788082379 1811928593 -16 23021 +M 6 1284263923 27 481598971 +M -1021604028 317026836 -20 -144 +M 1484710271 1291379399 -18 3406 +M 1531007314 1717284272 -15 37363 +M 62773 230477732 12 27595098 +M 1149411 1989816280 10 1090582237 +M 1231 1222749949 20 734963470 +M -2836807 1282718530 7 -216890866 +M 1249834025 663388522 -3 48261459 +M -208437 586567787 12 -233197078 +M -939062148 441089456 -17 -1472 +M 14 748292996 27 654756372 +M -84457 613961905 14 -395609898 +M -656488984 1858494846 -23 -68 +M 1984383787 660401369 -15 18623 +M 2 287573263 29 143786632 +M 1 1232008815 30 616004408 +M -15282594 1955722135 -26 0 +M 829533335 806346560 -12 76044 +M 1684 1778234474 20 1462181081 +M 751836801 1491087141 -31 0 +M 15894 143490430 16 69599515 +M -1687782717 746144887 -4 -36651340 +M -1598328112 2146252999 -1 -798706083 +M 625837553 1477293175 -7 3363477 +M -11452 103735904 17 -72508763 +M 259228643 1705099544 -8 804013 +M 1699804703 1063210770 -24 50 +M 230071112 1675973471 -30 0 +M 30038816 2022114299 -7 220978 +M 2131680203 438963828 -24 26 +M 0 114229663 28 0 +M -1526336731 1072290175 -15 -23259 +M -1820937242 1626893861 -18 -5262 +M -1785100930 326636178 0 -271517106 +M 1162 1732722627 20 983117037 +M -232587 359532766 10 -39874386 +M 154777973 1309076803 1 188701091 +M 2066242525 1469896163 -26 21 +M 1743473865 1037511163 -23 100 +M -927901229 1676075474 -8 -2828951 +M -8 1364622437 27 -682311218 +M 7004738 364058344 6 75999895 +M -85 705730463 20 -29290571 +M -1031131059 890757912 -2 -106926094 +M -1751483878 1912332889 -24 -93 +M -1486153598 286970545 0 -198596301 +M -1425679 554917969 10 -377242515 +M -15539589 1705911615 7 -1580069385 +M -15108 132097208 17 -121809364 +M -244005365 923959733 1 -209967729 +M -1379516510 1463967172 -7 -7347142 +M -3 1642190806 25 -76977694 +M -1875 2011200026 16 -115081789 +M -12412781 1992890926 4 -184307386 +M -103 1340610370 24 -1078772407 +M -1106237619 1731488046 -17 -6805 +M -54832622 785835874 3 -160520678 +M -1736763361 1907286139 -6 -24101644 +M 483123150 1344701147 -27 2 +M 952772487 440450248 -29 0 +M 21478097 1161473788 5 371728043 +M -10379024 223990148 7 -138568826 +M -2032600 1758444011 10 -1704317711 +M -567858006 987074620 -30 0 +M 5776 392598601 18 276812685 +M -1 1419545340 29 -354886335 +M -53 1340248642 23 -277473352 +M 57146 376754716 15 328522110 +M -1264257478 1683435960 -17 -7561 +M -115164 58511982 13 -25705238 +M -1901626963 1044336832 -8 -3612402 +M 1308637663 2115790781 -9 2518212 +M -1969362315 330689009 -24 -18 +M -153603 1875403820 13 -1098890888 +M 1090103428 435164070 -17 1685 +M 55692982 1513118476 -30 0 +M 245437 235704530 12 110341287 +M 231154 1084894792 13 956641276 +M -5993012 495212810 8 -353791274 +M 1978614125 1085626683 -27 7 +M 842767867 1914102748 -13 91697 +M 14532 167283161 17 148373956 +M 448412159 1196706219 -20 238 +M 12424 113898737 14 10796188 +M -2007516839 1584259461 -2 -370250496 +M 1 284316150 29 71079038 +M -303176240 1204805626 -12 -41526 +M -1458037960 246615885 -6 -2616255 +M -125 324664688 23 -158527680 +M 238546507 1426850018 -7 1238259 +M -12815 1413918646 17 -1105918423 +M -2136887886 545703360 -17 -4143 +M -194239 1119712431 13 -829665462 +M -5 447499283 28 -279687052 +M 1179 495010098 20 284969192 +M 499689789 1474196777 -31 0 +M -1770098 139258143 8 -29385157 +M -12 1302964444 27 -977223333 +M -56662 1215648748 15 -1051042013 +M -1652575912 1329531256 -22 -244 +M 370 101508240 21 36677782 +M -1186378889 381386002 -9 -411518 +M -1081332933 1915427023 -2 -241121082 +M 226103947 1720903969 3 1449522301 +M -1539156784 1283134200 -14 -56131 +M -13567 1478415064 17 -1224222240 +M 557278889 95356646 -6 386646 +M 2129749402 340625429 -29 1 +M 1875687180 1133911718 -13 120898 +M -5138 1037483365 17 -325353365 +M -2002670465 323116132 -3 -37665894 +M -2022612 805394562 10 -776768067 +M -68 63584093 24 -33779049 +M 15742148 388813498 6 182412852 +M -917376295 422185110 -12 -44031 +M -17650379 6694092 -2 -13755 +M -170172219 1422007061 -27 -1 +M 19 336909294 26 200039893 +M 28988058 629836667 -31 0 +M -9245038 2137475766 7 -1177850049 +M -178503119 723732655 -14 -3672 +M 358513354 1405235295 -24 14 +M 207768202 701921804 3 543285277 +M -554551108 1779742233 -31 0 +M -382915212 1215397017 -18 -827 +M -54 580187421 24 -244766568 +M -820559683 121613386 -3 -5808603 +M 167878418 349183333 -17 208 +M -1470847670 499221718 -5 -10685167 +M 1988549 1962880102 10 1861230499 +M 86733773 786700530 1 63547404 +M -176711 1481060186 13 -998381144 +M -1649593559 1731283550 -6 -20779511 +M 2015655126 1520019491 0 1426709387 +M -744833298 948966354 -3 -41142440 +M 839409995 1746057095 -6 10664064 +M -39 643911956 24 -196191924 +M -2134566853 1329182950 -18 -5040 +M -1887773 1708959315 7 -192292169 +M 9 65079318 24 4575890 +M 60987 187185851 15 174192863 +M -1324316879 1111895278 -12 -167404 +M -1241187066 468364652 -30 0 +M 122069558 987296993 -5 1753781 +M -40 532425299 25 -332765812 +M 1108761328 1650321222 -7 6656819 +M 1959378620 2044070106 -13 227664 +M 1450870583 517725379 -10 341585 +M 3 189633840 29 142225380 +M 474465 1376103319 12 1245332453 +M 57 251172239 23 55925069 +M -1553251961 1605698636 -4 -72586541 +M 9 1814104685 27 1020433885 +M 1916768993 562175428 -1 250889089 +M -2060977031 1804654477 -27 -13 +M 1 269527845 30 134763923 +M -144158487 1705752766 -4 -7156595 +M 223941001 1243416834 -16 1979 +M 299189535 1560103422 -9 424522 +M 1772458709 545808332 -28 2 +M -42272 1939040265 15 -1250718843 +M -1889706978 427415542 -30 0 +M -583197 1674953887 11 -931575853 +M 46910 546305215 14 195519849 +M 2044622470 2129350311 -2 506839423 +M 150618 1225213574 13 703961251 +M 5786 936621635 16 165383691 +M 1236906 1073109905 10 632923164 +M 459 2110450413 21 945992910 +M 32762 405855485 16 405781171 +M 791360021 1841308489 -25 20 +M 47388511 1104111385 5 779661454 +M -968703959 1981702229 -5 -27935061 +M -439140018 1848839147 -29 -1 +M -871463434 760364009 1 -617121747 +M 260530414 2144595923 2 1040720313 +M 14074787 896324967 6 375973671 +M -1 1055325461 30 -527662730 +M -1641221793 915093145 -8 -2731887 +M -180067869 1750162321 -1 -73376112 +M 1696222498 806188605 -22 152 +M 3419 1159605734 18 483971192 +M 210927 1130433260 12 454786103 +M -249584922 1044779261 0 -121426373 +M 71511758 883653416 -4 1839118 +M 1135201730 74509694 -23 5 +M -5 1335872397 27 -417460124 +M 93615 2131235270 14 1522183150 +M 2 1683697111 29 841848556 +M -340 2113163250 22 -1403272471 +M 545859911 494911784 -26 2 +M -60644503 458920376 4 -207357095 +M 1122541824 508012589 -21 127 +M 43584961 1244074332 5 807984639 +M -1792465069 428869842 -5 -11186555 +M -3 441406374 29 -331054780 +M 65795629 202753119 5 198785499 +M -7576 413437422 17 -191174433 +M 986443 1008804598 9 237257060 +M -1832936562 1264399221 -4 -67449977 +M 1896683631 2081264554 -19 3506 +M -970491634 699939293 -1 -158158417 +M -1685282218 1402466516 -6 -17197103 +M 5 1622904910 28 1014315569 +M 2115150743 1520620454 -26 22 +M 21251678 507676509 5 160768296 +M 5139 1708280480 18 1071637376 +M -1651339575 1824320745 -22 -334 +M -494143504 352253145 -30 0 +M 2009257559 10874387 -4 635903 +M -400864349 1971291495 -27 -3 +M -1846672998 1432228010 -27 -9 +M 1131298407 858662471 -22 108 +M -35 171878997 25 -93996326 +M 232238 65510029 13 58036492 +M -24123284 1684548269 -17 -144 +M 884471124 2094300769 -23 103 +M -1974949610 1901199047 -19 -3335 +M 1656682933 1284071592 -22 236 +M -726229389 579931754 -7 -1532184 +M -503170 74721614 10 -17927968 +M 1125552 976690928 10 524194921 +M 1742 1923330709 20 1635958054 +M 690862561 799792617 -27 2 +M 1271345829 1000416109 -16 9037 +M 1412730326 331464889 -16 3327 +M -587732751 1651273110 -14 -27583 +M -1122366 1928154719 9 -515960526 +M 1081896447 1658598412 -20 797 +M 2122807213 1820467365 0 1799548628 +M 55752805 409863829 4 170253650 +M 1954 792674297 20 756291785 +M 809034363 70968772 -15 816 +M 462647533 345177827 -18 284 +M 172370223 656523462 -9 102923 +M -2027576576 17547451 0 -16567670 +M -21 415881939 26 -272922522 +M -4 1918841913 27 -479710478 +M 16711215 1638904266 7 1632456872 +M -901120579 150108227 -11 -30756 +M -101 1868394578 22 -368570024 +M 1405435816 1738942476 -7 8891118 +M 921193645 386252777 -28 1 +M -3644772 1552944454 9 -1349479786 +M 1877689689 980737214 -20 818 +M 46899384 1386922135 4 484628929 +M 969304 912361719 11 843387474 +M 1854794801 188604094 -20 155 +M 9 212271297 27 119402605 +M 1833255379 597261450 -28 2 +M -412582958 1167702223 -14 -13693 +M 1243879177 71679490 -24 2 +M 72785259 729862980 -27 0 +M 1976845758 1347666056 -28 5 +M 22748386 1862905935 5 631482948 +M -1145595616 500478238 -25 -8 +M -9528413 347665316 3 -12340764 +M -286587291 2098223632 -28 -1 +M 1634235360 57503105 -1 21879936 +M 144 372406387 20 26184824 +M -2 1414888698 28 -353722174 +M 44 861713291 25 592427888 +M -1972110689 1206636319 -3 -138512136 +M -1947268503 600146948 -25 -16 +M -1973692841 1203304449 -8 -4320015 +M -1208244690 1890573267 -31 0 +M -397084918 42051941 -26 0 +M -47682 2046577168 14 -744513645 +L -13126466337780 1619327080 -22 -2359898 +L 139428529009474 1778304108 -2 -1476305180 +L 34770726550097 2029006530 -19 62660646 +L -2962303578556 919024996 5 -814242884 +L -36017923099112 1803256649 -11 -1883179598 +L -70172277282120 106332436 -2 -1325343611 +L -42959967947305 1332204675 -25 -794253 +L 115693242742799 1237485449 -30 62091 +L -60698031136807 1608662519 -12 1784327512 +L 35584593287601 49561991 -17 6263599 +L -52426050562345 303332610 -14 -451929424 +L 26090926090811 332463043 -9 -700707437 +L -39759554905764 528942147 -10 -973611817 +L 30342881081093 2045925796 -2 -1531295918 +L -13407068078141 716805166 -15 -136575206 +L -113297134738158 347978351 -5 1787724121 +L -4055418436313 2098957044 -9 848069572 +L 111098396071469 233688945 -14 737936944 +L -70084317059393 1312802125 1 -135757152 +L -97015134147506 1510935838 -23 -8137009 +L 26879542040872 1752545686 -24 1307513 +L -87403422412212 496808167 -3 2096938413 +L -18177172973453 1806178832 0 1909527484 +L 136945646055211 880487520 -10 -1002332277 +L 93751418084425 985452883 -21 20514418 +L -24921807012857 1135928545 -20 -12571972 +L -57662037604331 917690258 -27 -183591 +L 120747870744859 937306585 -7 -583221797 +L -64773591462692 99817764 -28 -11215 +L 82791200409815 253875262 6 -2008797661 +L -23555224810115 154094981 6 -646257841 +L 57355897697212 1448439803 -15 1180565632 +L 137352439197362 2095662567 -28 499326 +L -77118644440602 1060991151 -31 -17742 +L -36953855653426 1895133931 -12 628284072 +L 130223186572762 1829869332 -15 -908592905 +L -101812934377576 128731988 -12 -1489822590 +L -64790041450827 1973901340 -25 -1774798 +L -62354160921342 1379363676 -25 -1193592 +L -24086925438749 1890971234 -29 -39506 +L 113290775550283 2614224 3 -1746332385 +L 113027163917907 762481489 -22 9568401 +L 9272825680919 868420431 -17 28608882 +L 138425114666506 1407797456 -13 -1807714349 +L -84958991748310 2132582225 -23 -10057740 +L -7875432461884 1005106542 4 1192242698 +L -140206447773985 1362289249 -11 -479335611 +L 86609734226049 426315644 -4 851378859 +L -127133161043471 1714658875 5 897998787 +L 71660115287244 775932240 -1 1376116868 +L 135025627814740 1532179850 4 -2001664193 +L 116695185820907 1446264050 -22 18737260 +L -11683590272085 404690250 -20 -2099730 +L 12997656736942 423214262 3 1598532535 +L -13029813795693 1456937008 1 -1711392626 +L -49782557429648 422529608 7 1546525397 +L -122101158406361 992633130 3 -188397363 +L 5254346181626 809922719 -15 60473764 +L 1659785624653 175125475 -1 -1047553075 +L -79842280066125 989745072 -5 1131617142 +L 19041645946254 1793618208 -31 7406 +L -124367836699283 71742402 7 2070341269 +L 41473295854075 1333747861 -22 6141076 +L -5778147223301 1009496979 -15 -82893837 +L 2294100807131 190460998 -31 95 +L 122752797169999 1343409465 -17 585873981 +L 95578380089564 1556510246 -1 -1665650090 +L -104049978640429 344159811 -9 1793760923 +L 92543094615692 788607488 -16 518546932 +L -8178440853449 999400541 4 -693385356 +L -24936217111797 214545062 7 -927493749 +L 75554037791694 961969572 6 -1277795269 +L 28559012862882 248077384 -6 4516902 +L -114019836564056 43012163 -11 -1114562344 +L 46531192734930 1327716857 -30 26792 +L -140261211437297 1627702678 -29 -198024 +L 73689106536429 1363622582 1 -935497049 +L 118950670795668 1836272975 -19 193999082 +L 85637168861755 1266057238 -1 -1303830948 +L 29514435135541 1004065802 -20 13160480 +L 38792518756999 253926590 -23 546865 +L 125444088650469 1631080661 -12 1786267138 +L 113155533079877 1974262677 -3 -1581130555 +L -75781818786894 753455195 -26 -396205 +L 39283245091392 354762846 -24 386790 +L -120557372143123 1989808032 1 490288089 +L -76003128917551 819742158 -21 -13833737 +L -31673645284631 677683393 -28 -37237 +L -68407924383900 49544925 -6 1109506938 +L -24482880398137 1614953645 7 1519327883 +L -101937904509263 793440408 -25 -1122464 +L 18526160996920 56676435 -15 14924565 +L -101730221563931 164579023 -30 -7260 +L -105011842030620 1295886648 -4 -657531921 +L -122131586537841 537871639 -13 560982428 +L -134720070740595 1410911160 -10 -538964315 +L 92483579279040 1752171519 -31 35138 +L -8539473532514 2072141685 -24 -491129 +L -48077046151944 1663091941 -15 -1136261225 +L 118893684994215 2086104930 -18 440574354 +L 51483645605455 957784120 0 1552361332 +L 123394861290216 587870018 -9 1549005223 +L -14908382041242 1034700215 -20 -6850271 +L 50064288734979 1736066027 -1 -1633936631 +L -31410628841653 1607243894 7 -1320824140 +L -125399350603451 1237773368 2 -2040046355 +L 11507354012307 1580454145 -4 1029915385 +L 52122887819481 833767564 -3 -181832469 +L 15228438208020 2018093526 -3 -2119960692 +L -99729744585927 310125218 -11 1557761444 +L 89102059379703 346904265 3 1434770035 +L 47464463877081 251618128 -20 5303186 +L -132243090863776 34090964 6 -1466845306 +L -79901217869055 613158545 -9 -1608124614 +L 1091094406748 697677081 -14 21636097 +L -64955717025243 904540042 -8 500812422 +L -208789661212 1562445836 -16 -2317947 +L -114630378716644 495103697 -19 -50409727 +L -133947971504710 10121693 -28 -2345 +L -104127208704336 889676492 -25 -1285595 +L -41629487892566 1459715598 -22 -6746367 +L -92085786138987 1267342059 -21 -25913395 +L -125527862875484 1712730074 -11 -1639292102 +L -135533637814142 1189881649 -12 -1154137544 +L 98928463569801 782552204 2 -1282885198 +L 87026374204564 771125089 -2 -409289417 +L 112568982884791 1985987113 -7 1566204798 +L -12364572991590 812290300 -15 -142733457 +L 114393927745781 362490509 -29 35966 +L -25387860186159 1305551164 -7 -321484686 +L 3053767040060 494845523 6 72386279 +L 136075136554950 1729508408 -11 1970802842 +L 49990184054748 1696115080 -21 18827209 +L 56955723978526 60780347 -21 768311 +L -136491523336563 179620575 -31 -5317 +L -32185140649168 1968081950 -23 -3516296 +L -21868719214226 1927294192 -19 -37434237 +L 35185875658642 777135943 -24 758944 +L 69466623456307 2001614577 -14 -343086935 +L 5185298956284 1441840465 -16 53123460 +L 102562902387345 279002062 -4 -455643448 +L 40811796141907 1510415477 4 -1537822870 +L -120246574784595 1354077052 -6 692494444 +L -94296266821038 420646277 -14 -1127436267 +L -69901603955285 1485278178 -4 1939996322 +L 58830717024207 703878970 7 -596261288 +L 38395205814315 1721015758 -24 1834081 +L -118411753326276 1995571955 -9 -164400750 +L 6539791772924 2134819803 6 -2068213109 +L 140454294959089 1057478914 -4 2021581445 +L 106662993636161 645644353 -17 244668641 +L -35562544095147 806698845 3 1711444345 +L 100747803988983 424863303 -6 -2087370079 +L 109588988705987 2084941667 2 1156254617 +L -104368902735818 587291441 -6 714771784 +L 19234930012739 1370740035 -11 1700035036 +L -39159594848933 1021688311 -28 -69406 +L 5501038877865 224031620 -7 187910169 +L 9963812740613 1906389872 1 -702420417 +L -54999192367546 2145453268 -14 941258760 +L -67479863000393 693954581 -1 1841683307 +L 83414129568893 1957971822 -28 283317 +L -114514950424878 2086104428 2 1311743261 +L 63897331089997 719305331 -5 -1167712999 +L -19007187116744 1758945756 -22 -3711713 +L -49938467747357 1969669557 -8 1467326480 +L 13978404011395 481508396 7 -2145783977 +L -36218496657079 1012641561 -12 125263199 +L 61964954075547 835585729 -30 22455 +L 98827868217618 1554223425 -29 133230 +L -60670503560475 962409013 -30 -25322 +L -68606597767999 144975900 -26 -69011 +L -80655778881119 2067646364 -27 -578595 +L 36196561663696 2018439435 -3 675181984 +L -117375800122445 1135853210 7 -1062815757 +L -4375682020836 1401821151 -5 934365987 +L -38106194802138 1080576281 -5 2106674132 +L -84022559251344 1000416833 -8 1720467295 +L 109268323274063 1312136086 -15 2037519932 +L 11070103801503 1090824985 2 161272803 +L -29818149802015 1966152365 6 -1282686368 +L -89812752311246 750387954 0 391035419 +L -140260974501193 2044228947 -12 1763296886 +L 98349227367728 1240874675 -10 -338350952 +L -65570155392918 1255514092 -13 -384716438 +L 70781907168487 1620330938 -15 1629825563 +L -83330500002397 2109953794 -23 -9760062 +L 95397145085967 2042699224 -23 10817288 +L -128921073260775 97498105 -19 -11166241 +L -108669300112071 1247229929 -6 1702488974 +L 108636619160344 1402048522 -4 599713503 +L -3873477168812 1432699341 -6 -1722951457 +L 73841728825188 211953446 -4 216237583 +L 20348503073976 1934578019 -29 34144 +L 85790432860715 2082667814 -5 1579657861 +L 131274095320326 1968056194 -12 -693382201 +L -62008863818001 1563243415 -19 -86094802 +L -84794559178004 584974262 -14 -1409791848 +L -84488274050902 107647235 6 -1740740062 +L 10333609916166 1599456900 -22 1835012 +L -67572802831451 343870236 -11 -988306408 +L -69417923303530 922854082 4 881821169 +L -88753947884544 2080522489 -22 -20500613 +L 3028585130723 224527898 -28 1180 +L -93628760136578 1991290037 -20 -82797775 +L 130201485688493 1724221397 -29 194723 +L -94177398829034 771202852 -31 -15750 +L -60568294475670 1480811093 -12 -1606487550 +L 22280353764367 1411898659 -8 1386830837 +L -119647138122490 271210535 -5 282384357 +L 12451497238229 1015288759 -27 43860 +L 23330223548091 547561162 -30 5540 +L 87842360180188 965280770 -12 1049851019 +L 98037301666754 895424248 2 1783111697 +L -75230436530054 1826120905 -17 -488064458 +L -125420892881134 1635793494 -21 -45554850 +L 27614647055608 752493004 -1 1992091618 +L -115256521392316 1074195943 -15 -1759426335 +L 84436125852484 2128763725 -13 1627248640 +L -112223514647855 528844868 -29 -51480 +L -16151613722099 307209689 -15 -70518595 +L 17040995985901 1651463114 -27 97638 +L -1616250358778 1252115230 -8 613774975 +L -116688608307776 819528985 -10 -537609045 +L 110742421838816 423293186 -29 40659 +L 108116781725871 168277859 -28 31564 +L -5842486338721 1483180343 -21 -1924158 +L 61987331683499 34841505 -12 245699737 +L 50947015440531 1782196152 -19 80643986 +L 31610368850692 272878255 -13 490343481 +L 43270865376908 78899962 -2 -1954832110 +L 115170370052071 95464151 -18 19534867 +L 30334053269316 256007453 -17 27586895 +L 20050764670440 1115182843 -12 -1752949809 +L -92997078696626 1797271569 1 1373843677 +L -73546533989855 1962774066 -9 1851720940 +L 22436598787409 1845291584 -25 574571 +L 12665354525336 1984967163 -13 1429052121 +L -75312015570332 1328911873 6 1208555783 +L -80430954601687 1374199797 3 1510994968 +L 94351928616658 1735862118 -27 568229 +L -71583785304950 1420175833 -21 -22573231 +L -13970495651039 902716976 -16 -89607019 +L 125119930707618 57449829 -22 798392 +L -76886431335553 421687476 -13 -1842853797 +L -133135289331194 21660462 4 315074530 +L -5431644371200 1773919675 6 1457037714 +L 116270048660852 335263516 -1 1235094594 +L 105993830071014 755817092 -29 69487 +L 114703233678065 201509376 -6 682674706 +L -136637058111717 459410189 -31 -13612 +L -105664826461816 1930329945 -28 -353822 +L -109021053625386 106548807 -21 -2579592 +L -16938346481128 702811344 1 -1531917790 +L -15730331342175 1379421527 -30 -9410 +L -55098416824617 311561293 -31 -3722 +L 90255141413780 1749041732 3 1733605840 +L -19780741734485 159282352 3 -1306186878 +L 105227114987526 1776526556 -9 -1776643302 +L 134313464504757 1127737592 2 573081358 +L 58606002035977 1439055408 -23 4681608 +L -135785044159249 818659674 -13 -2023969840 +L -85537437418676 940927714 -8 -367389962 +L -30660444736032 1892346968 -27 -201299 +L -71889356189743 1921105408 6 -1990492275 +L -15065823580072 1020894794 -24 -426908 +L 63095523119755 1046863836 -11 2133797971 +L 15873070857295 2017848259 -5 -2060420692 +L 60391364490528 329750603 -1 -1581248950 +L 105331532600622 178299906 -9 -96754786 +L 49845415357403 1988967355 -25 1375846 +L 75628099796417 1607775613 -13 -1678089935 +L 29737765910294 2012917847 4 187702920 +L -33840798789982 404225558 -9 443621586 +L -6671705973866 448469463 -22 -332180 +L -25005224910848 262957075 0 758876191 +L 12703501337611 400563803 -15 72311424 +L -63969518073654 1547205285 3 -2139620600 +L 66728590184655 1260240618 0 -1636644683 +L -77549053336775 1714447284 0 1454447999 +L 9040734939049 1090930616 -9 380092008 +L -38497061254291 228743667 -15 -125127606 +L 107421002064374 729615527 -16 556892723 +L -91893464525277 651950348 -5 71086177 +L 115301415804448 1308202321 -4 586419491 +L 7957187165041 1425538146 -10 863359783 +L 132108662766623 1560673628 -16 1464987032 +L 29663180492259 1459883454 -2 -963109368 +L 98333660397290 240874750 -3 -143692537 +L 82749163233534 585872654 -21 10765180 +L -122720434935964 2141152927 4 296493106 +L 37435459380652 889893287 -14 946849776 +L 121313892191547 1707559389 0 311586867 +L 91920449606629 1748610123 -20 71380678 +L 51339663957352 1794468355 5 1293611453 +L -60630103302569 228180798 -10 -1996720780 +L -7921417117813 607328627 -22 -534112 +L 67083289307349 2009440468 -19 119727793 +L 128093146825128 764809448 -7 -82982263 +L 67818784980179 560170276 -10 96983364 +Q 0 0 0 +Q 1 1073741824 1 +Q 0.5 1073741824 0 +Q 0.25 1073741824 -1 +Q 3.7000000000000002 1986422374 2 +Q 1e-10 0 0 +Q 4.6566128730773926e-10 1073741824 -30 +Q 2.3283064365386963e-10 1073741824 -31 +Q 0.99999999900000003 2147483646 0 +Q 1.5 1610612736 1 +Q 123.456 2071247978 7 +Q 0.0078125 1073741824 -6 +Q 0.75 1610612736 0 +Q 3.4840877667667769 1870505377 2 +Q 5.5063814203472754 1478108007 3 +Q 5.3610169876629241 1439087040 3 +Q 2.633305222586185 1413744976 2 +Q 7.1431816951165779 1917483236 3 +Q 0.41450280243242688 1780275981 -1 +Q 6.257463295989151 1679725013 3 +Q 1.3895329338067191 1491999627 1 +Q 1.2889174822309482 1383964608 1 +Q 3.7522824039128375 2014491276 2 +Q 3.5078412220085058 1883257916 2 +Q 5.6964929204778469 1529140675 3 +Q 0.94896064436881566 2037877466 0 +Q 7.3645415906080727 1976904080 3 +Q 0.013459589006783013 1849871827 -6 +Q 6.4961639062790146 1743800720 3 +Q 3.7766802549349405 2027589773 2 +Q 0.63226221488597822 1357772768 0 +Q 6.6395002772348457 1782277285 3 +Q 7.6271176570074086 2047388806 3 +Q 0.91957259036692562 1974767101 0 +Q 5.7034445447213082 1531006737 3 +Q 7.4441240635158685 1998266838 3 +Q 4.2806015374497219 1149065226 3 +Q 5.2780924588493194 1416827156 3 +Q 6.7308533266859092 1806799682 3 +Q 1.0603809244521978 1138575348 1 +Q 4.1986638607786837 1127070248 3 +Q 2.3513195109472953 1262355050 2 +Q 4.8327703738109822 1297286919 3 +Q 1.7956827659840693 1928099688 1 +Q 5.4723404882330806 1468970214 3 +Q 6.8001185123181136 1825392914 3 +Q 0.25216586608786162 1083044148 -1 +Q 3.0322850706680073 1627945651 2 +Q 2.3717878449545697 1273343903 2 +Q 7.3936414667070922 1984715519 3 +Q 4.1710287193965589 1119651996 3 +Q 0.22502924551628567 1932986500 -2 +Q 7.1495748862380308 1919199395 3 +Q 3.3266765942005687 1785995897 2 +Q 2.3581676029021192 1266031592 2 +Q 4.3288220553597263 1162009322 3 +Q 2.2371641547004275 1201068360 2 +Q 0.42698728881844872 1833896441 -1 +Q 5.7730171007554327 1549682478 3 +Q 6.3513761881572703 1704934563 3 +Q 0.98189587886310825 2108605344 0 +Q 0.16006166318429216 1374919217 -2 +Q 3.6023633000727422 1934004070 2 +Q 7.4606181089427945 2002694424 3 +Q 2.0471156902767378 1099036868 2 +Q 3.7763961468303 2027437243 2 +Q 1.636883812528203 1757590611 1 +Q 5.0928716895479074 1367107334 3 +Q 3.5083902788567825 1883552689 2 +Q 5.8237258223122454 1563294497 3 +Q 7.3400233585865244 1970322517 3 +Q 6.7687080240702295 1816961225 3 +Q 7.8224579976068105 2099825080 3 +Q 1.8969556043726863 2036840571 1 +Q 1.9994889456834679 2146934908 1 +Q 0.72129385002175805 1548966748 0 +Q 0.69842417058507067 1499854486 0 +Q 5.8594172835346647 1572875350 3 +Q 4.2698793646534963 1146187014 3 +Q 5.9679660608631373 1602013691 3 +Q 1.60594348124945 1724368683 1 +Q 6.468518834377333 1736379803 3 +Q 5.2773117194505934 1416617578 3 +Q 2.2873744720203524 1228024819 2 +Q 1.1471715017458424 1231766021 1 +Q 2.391864738997441 1284122604 2 +Q 7.1468991545060421 1918481134 3 +Q 7.8302313903346796 2101911734 3 +Q 3.982178649708267 2137915883 2 +Q 2.3753883359870844 1275276902 2 +Q 7.1480750867543632 1918796795 3 +Q 7.1548156438771544 1920606200 3 +Q 0.29579912040177714 1270447548 -1 +Q 1.8728048284511365 2010908872 1 +Q 5.1399448902452942 1379743450 3 +Q 2.4616629339716667 1321595224 2 +Q 3.0306911940071437 1627089945 2 +Q 6.7812332529909156 1820323441 3 +Q 1.1911549542190611 1278992893 1 +Q 4.2728970569889757 1146997070 3 +Q 1.796893744450486 1929399967 1 +Q 6.2795503943279485 1685653974 3 +Q 3.3445197419727428 1795575364 2 +Q 2.3913353074714463 1283838367 2 +Q 3.8440139174261416 2063739258 2 +Q 6.0230345997234389 1616796039 3 +Q 7.9639196538769728 2137798404 3 +Q 6.3303406034761585 1699287867 3 +Q 7.2222240419401942 1938701004 3 +Q 6.9265335809793616 1859327200 3 +Q 3.2836156871133078 1762877749 2 +Q 2.6790530988516386 1438305680 2 +Q 7.4202810866480062 1991866537 3 +Q 7.8752495818419606 2113996213 3 +Q 2.0095727229081639 1078881140 2 +Q 6.714815760149528 1802494631 3 +Q 5.1852461459076684 1391903914 3 +Q 0.73766249917297078 1584118155 0 +Q 0.93690198919707957 2011981702 0 +Q 7.4770281358125663 2007099457 3 +Q 2.6105237888315713 1401514287 2 +Q 6.6409972160882207 1782679116 3 +Q 0.48373292635830722 2077617099 -1 +Q 7.9628489792523505 2137510997 3 +Q 2.6457350040535639 1420418165 2 +Q 2.9738632366753395 1596580668 2 +Q 4.9413277561153759 1326427569 3 +Q 4.1850187087889186 1123407405 3 +Q 0.50022280053440671 1074220285 0 +Q 3.434136375539131 1843687928 2 +Q 3.8027093019031026 2041564011 2 +Q 6.9419745225758964 1863472097 3 +Q 6.2782150899236573 1685295531 3 +Q 1.3730070169601394 1474255059 1 +Q 3.2110165822683898 1723901401 2 +Q 6.1423418665755243 1648822340 3 +Q 6.2354329083759836 1673811276 3 +Q 1.9502886384747722 2094106480 1 +Q 6.1790404358239899 1658673537 3 +Q 1.817185459071319 1951188029 1 +Q 2.9690271892164244 1593984335 2 +Q 2.699118962447784 1449078459 2 +Q 1.2399900261744978 1331429152 1 +Q 2.586091408266733 1388397253 2 +Q 7.1775908031094042 1926719860 3 +Q 3.2467010063893511 1743059330 2 +Q 6.9295524406262068 1860137569 3 +Q 6.7487533134178994 1811604673 3 +Q 7.253481762706218 1947091685 3 +Q 3.4985693794566624 1878280133 2 +Q 2.939882600370292 1578337453 2 +Q 0.35665441414284782 1531819045 -1 +Q 3.3451873582213834 1795933788 2 +Q 2.5571164507096373 1372841441 2 +Q 4.9633797991133388 1332347120 3 +Q 7.9259937393950661 2127617744 3 +Q 7.1101073008918814 1908604896 3 +Q 7.5136527615444786 2016930805 3 +Q 2.0811047967463328 1117284630 2 +Q 7.0499960233705314 1892468897 3 +Q 7.0380702915812074 1889267608 3 +Q 7.1332684627691556 1914822173 3 +Q 0.47440081495140252 2037535985 -1 diff --git a/tests/data/tflm_requant_oracle.cc b/tests/data/tflm_requant_oracle.cc new file mode 100644 index 00000000..fa3d82b3 --- /dev/null +++ b/tests/data/tflm_requant_oracle.cc @@ -0,0 +1,38 @@ +// Generator of tests/data/tflm_requant_golden.txt. +// Built against the real TFLite-micro sources (AdaWorldAPI/tflite-micro @ 2d8adcc, +// default build: TFLITE_SINGLE_ROUNDING undefined = double rounding): +// g++ -std=c++17 -O1 -I tflm_requant_oracle.cc \ +// /tensorflow/lite/micro/kernels/internal/common.cc \ +// /tensorflow/lite/micro/kernels/internal/quantization_util.cc +// ./a.out > tflm_requant_golden.txt +// Not built by cargo; kept so the fixture can be regenerated. +#include +#include +#include +#include "tensorflow/lite/micro/kernels/internal/common.h" +#include "tensorflow/lite/micro/kernels/internal/quantization_util.h" +using namespace tflite; using namespace tflite::micro; +static uint64_t s = 0x1200000000000019ull; +static uint64_t nx(){ s += 0x9E3779B97F4A7C15ull; uint64_t z=s; z=(z^(z>>30))*0xBF58476D1CE4E5B9ull; z=(z^(z>>27))*0x94D049BB133111EBull; return z^(z>>31);} +int main(){ + const int32_t E[] = {INT32_MIN, INT32_MIN+1, -(1<<30), -2,-1,0,1,2,(1<<30), INT32_MAX-1, INT32_MAX}; + // srdhm + for (int32_t a: E) for (int32_t b: E) printf("S %d %d %d\n", a,b, SaturatingRoundingDoublingHighMul(a,b)); + for (int i=0;i<400;i++){ int32_t a=(int32_t)nx(), b=(int32_t)nx(); printf("S %d %d %d\n",a,b,SaturatingRoundingDoublingHighMul(a,b)); } + for (int i=0;i<200;i++){ int16_t a=(int16_t)nx(), b=(int16_t)nx(); printf("H %d %d %d\n",a,b,(int)SaturatingRoundingDoublingHighMul(a,b)); } + // rdpot incl ties + for (int e=0;e<=31;e++){ for (int32_t x: E) printf("R %d %d %d\n",x,e,RoundingDivideByPOT(x,e)); + for (int i=0;i<8;i++){ int32_t x=(int32_t)nx(); if (e>0 && (i&1)) { x = (x & ~((int32_t)((1ll<0?shift:0; + int32_t x = (int32_t)nx(); x = (int32_t)((int64_t)x >> left); + printf("M %d %d %d %d\n", x,m,shift, MultiplyByQuantizedMultiplier(x,m,shift)); } + // int64 overload + for (int i=0;i<300;i++){ int shift=(int)(nx()%39)-31; int32_t m=(int32_t)(nx()&0x7FFFFFFF); int64_t x=(int64_t)(nx()%(1ull<<48)) - (1ll<<47); + int64_t red = m < 0x7FFF0000 ? ((m + (1 << 15)) >> 16) : 0x7FFF; (void)red; + printf("L %lld %d %d %d\n",(long long)x,m,shift, MultiplyByQuantizedMultiplier(x,m,shift)); } + // quantize_multiplier + const double D[] = {0.0, 1.0, 0.5, 0.25, 3.7, 1e-10, 4.656612873077393e-10, 2.3283064365386963e-10, 0.999999999, 1.5, 123.456, 0.0078125, 0.75}; + for (double d: D){ int32_t q; int sh; QuantizeMultiplier(d,&q,&sh); printf("Q %.17g %d %d\n", d,q,sh); } + for (int i=0;i<150;i++){ double d = (double)(nx()>>11) / 9007199254740992.0 * 8.0; int32_t q; int sh; QuantizeMultiplier(d,&q,&sh); printf("Q %.17g %d %d\n", d,q,sh); } +} diff --git a/tests/tflm_requant_golden.rs b/tests/tflm_requant_golden.rs new file mode 100644 index 00000000..af935e2f --- /dev/null +++ b/tests/tflm_requant_golden.rs @@ -0,0 +1,96 @@ +//! Byte parity of the fixed-point requantization chain against TFLite-micro. +//! +//! `tests/data/tflm_requant_golden.txt` was produced by +//! `tests/data/tflm_requant_oracle.cc`, compiled against the real TFLM +//! `fixedpoint.h`, `common.cc` and `quantization_util.cc` (default build, +//! double rounding). Each line is ` `. +#![cfg(feature = "std")] + +use ndarray::simd::{ + multiply_by_quantized_multiplier_i32, multiply_by_quantized_multiplier_i64, rounding_divide_by_pot_i32, srdhm_i16, + srdhm_i32, I32x16, +}; + +const GOLDEN: &str = include_str!("data/tflm_requant_golden.txt"); + +fn rows(tag: &str) -> Vec> { + let v: Vec> = GOLDEN + .lines() + .filter(|l| l.split(' ').next() == Some(tag)) + .map(|l| { + l.split(' ') + .skip(1) + .map(|t| t.parse::().unwrap()) + .collect() + }) + .collect(); + assert!(!v.is_empty(), "no golden rows for {tag}"); + v +} + +#[test] +fn scalar_chain_matches_tflm() { + for r in rows("S") { + assert_eq!(srdhm_i32(r[0] as i32, r[1] as i32) as i64, r[2], "srdhm_i32 {r:?}"); + } + for r in rows("H") { + assert_eq!(srdhm_i16(r[0] as i16, r[1] as i16) as i64, r[2], "srdhm_i16 {r:?}"); + } + for r in rows("R") { + assert_eq!(rounding_divide_by_pot_i32(r[0] as i32, r[1] as i32) as i64, r[2], "rdpot {r:?}"); + } + for r in rows("M") { + assert_eq!( + multiply_by_quantized_multiplier_i32(r[0] as i32, r[1] as i32, r[2] as i32) as i64, + r[3], + "mbqm {r:?}" + ); + } + for r in rows("L") { + assert_eq!(multiply_by_quantized_multiplier_i64(r[0], r[1] as i32, r[2] as i32) as i64, r[3], "mbqm64 {r:?}"); + } +} + +#[test] +fn quantize_multiplier_matches_tflm() { + for l in GOLDEN.lines().filter(|l| l.starts_with("Q ")) { + let t: Vec<&str> = l.split(' ').collect(); + let d: f64 = t[1].parse().unwrap(); + let want = (t[2].parse::().unwrap(), t[3].parse::().unwrap()); + assert_eq!(ndarray::simd::quantize_multiplier(d), want, "quantize_multiplier({d})"); + } +} + +/// The vector methods on whatever backend this build selected, 16 golden rows per call. +#[test] +fn vector_methods_match_tflm() { + fn chunks(v: &[Vec], f: impl Fn(&[Vec]) -> ([i32; 16], [i32; 16])) { + for c in v.chunks(16) { + let (got, want) = f(c); + assert_eq!(&got[..c.len()], &want[..c.len()], "rows {c:?}"); + } + } + let lane = |c: &[Vec], k: usize| -> [i32; 16] { + let mut a = [0i32; 16]; + for (i, r) in c.iter().enumerate() { + a[i] = r[k] as i32; + } + a + }; + let s = rows("S"); + chunks(&s, |c| { + let g = I32x16::from_array(lane(c, 0)).saturating_rounding_doubling_high_mul(I32x16::from_array(lane(c, 1))); + (g.to_array(), lane(c, 2)) + }); + let r = rows("R"); + chunks(&r, |c| { + let g = I32x16::from_array(lane(c, 0)).rounding_divide_by_pot(I32x16::from_array(lane(c, 1))); + (g.to_array(), lane(c, 2)) + }); + let m = rows("M"); + chunks(&m, |c| { + let g = I32x16::from_array(lane(c, 0)) + .multiply_by_quantized_multiplier(I32x16::from_array(lane(c, 1)), I32x16::from_array(lane(c, 2))); + (g.to_array(), lane(c, 3)) + }); +}