代码

/* References:
 * saffah_codex_6a_agg3 https://duck.ac/submission/125916: retained the accepted million-character suffix engine, input scanner, and inherited source citations.
 * Idea: Store decimal leading groups in32bits instead of64bits while retaining
 * the independent width table. All groups have at most four digits. Their
 * trailing space beyond byte3, when present, is redundant because the three-
 * digit suffix overwrites that position with a digit whose space bit is set.
 * This reduces the active prefix table from8KiB to4KiB without adding width
 * decoding, special paths, or output arithmetic.
 * Purpose: Measure the effect of a smaller hot decimal-prefix table.
 */
/* References:
 * saffah_cc_v41_agg1, https://duck.ac/submission/122408: retained accepted pinned LCP/table pointers, packed suffix engine and inherited citations.
 * Idea: In the token scanner, a block containing only signed bytes greater
 * than32 cannot contain any original delimiter. Accept such blocks with one
 * AVX2 compare. At the first other byte, resume the original exact scalar
 * delimiter test; unusual control/high-bit bytes therefore retain their
 * existing meaning. The copy, suffix ordering and output paths are unchanged.
 * Purpose: Official experiment reducing the common lowercase token scan cost.
 */
// ===== n06k_ 席 发件 [N06K-LPPIN]：walk 快环 LCPB 的 4 字节存数改**指针形态**并钉进寄存器 =====
// 基座 = 本文件正文 = problems/1006e6/work/n06i_v_nolp.cpp（= 本账号 #122204，7.944293 ms，榜 1）。
// 【本发唯一变量】walk 快环（`while (rp < rend)`）里给 LCPB 存 4 个 LCP 字节那一条：
//   旧：`*(uint32_t *)(LCPB + (int)(rp - RB)) = w32;`
//       ⇒ JUDGE 构建（g++-9 -O2，**不加 -DBENCH**）每迭代 6 条：
//          `mov rsi,rdx` / `sub rsi,r15` / `sar rsi,3` / `movsx r8,esi` /
//          `lea r12,_ZL4LCPB[rip]` / `mov DWORD PTR [r12+r8],ecx`
//   新：外层每轮进入快环前建局部指针 `uint32_t *lp = (uint32_t *)(LCPB + x);`
//       环内 `*lp = w32; lp++;`，环尾屏障改 `__asm__ volatile("" : "+r"(tb), "+r"(lp));`
//       ⇒ JUDGE 构建每迭代 2 条：`mov DWORD PTR [rcx], r9d` / `add rcx, 4`。
// 【与 [N06I-NOLP] 的关系（不是反向）】NOLP 之所以删掉跨回边的 `lp` 活值，是因为
//   gcc **把它落进栈槽**（JUDGE 构建 `add QWORD PTR 88[rsp],4` + 载回 = 每轮 store+load 链）。
//   本发把同一个指针**用空 asm 屏障钉进寄存器**（落 `rcx`），于是既拿回指针形态，
//   又没有栈溢出链 —— 即 [N06E-BASEPIN] 那把尺子（每迭代少 1 条 ≈ 0.1~0.4% 全时）用在 LCPB 基址上。
// 【JUDGE 构建读数（work/n06k_cen.py，即 `sh ref/gcc9/g9.sh -O2 -static -U_FORTIFY_SOURCE -std=c++17 -S`）】
//   walk 内环 **46 → 39 条**、栈引用 0（前后皆 0）、`lea _ZL4LCPB[rip]` **1 → 0**；
//   MSD 环 77 条 / 栈引用 1（**不变**）；函数总条数 2834 → 2831。
// 【本地闸门（判题机同款 ref/gcc9/g++-9，`-O2 -static -U_FORTIFY_SOURCE -std=c++17`，work/n06i_gate.py）】
//   (1) 6 档规模 n=1e6/2e5/5e4/1e4/7/3：fb/olen/ohash 与基座**逐位相同**；
//   (2) wrap 直跑 141 例**逐字节相同**；(3) 暴力 oracle 300 例（SA 与 LCP 全对）。
// 【本发思路 概览 / 预期】walk 每组 4 元素的机码少 7 条；按本题已标定的 walk 边际率
//   （[N06E-BASEPIN]：每迭代 6 条 ≈ 60 µs）预计 −60~140 µs 量级；目标只差 **7.24 µs（0.091%）**。
// ===== 以下为基座正文（引用区 → 思路区 → 代码），一字未动 =====
// ===== n06i_ 席 候选刀 [N06I-NOLP]：去掉跨回边的 lp 活值（基座 combo1 = #122121）=====
// ===== n06i_ 席 本发（第 2 代；正文 = n06i_sub1.cpp（本账号 #122038，8.022221 ms）+ 本席两把刀）=====
// 【本发改动一 [N06I-NTL]】LCP 行输出改**非临时存**（未对齐先普通存、之后每轮 +32 恒对齐），收尾 sfence。
//   SA 行**不做**非临时存（实测 8B movnti 因零散写填不满 64B 行反而 +6.9%）。
//   判题机相位 probe（`tools/probe.py` + 本件 -DBENCH -DPHASE，n=1e6，各 2 reps）：
//   afterwalk 962K → **731K**（−231K），total 28,404,934 → **28,163,130（−0.85%）**，其余相位不动。
// 【本发改动二 [N06I-MARK]】walk 快环四条守卫合并：把 D5B32 表里 q = lzcnt(记录异或) >= 44 的字节
//   打 0x80 标记（键相等 <=> 异或 < 2^20 <=> q>=44；LCP 字节值域 0..6 永不含 0x80），
//   四条 `cmp+jbe` 换成对 w32 的一次 `test 0x80808080 + jnz`。断点仍在四次发射之前，逐位等价。
//   gcc9 快环 96 → 87 条指令；probe：lcpline 8324K → 8250K，total −0.15%。
// 【本地闸门】6 档规模逐位相同 + 141 例 wrap 逐字节相同 + 300 例暴力 oracle 全对（`work/n06i_gate.py`）。
// 【判题机读数】基座 #122038 = 8.022221 ms。
// ===== 本发思路 结束 =====
// ===== REFERENCES（引用区）: n06i_ 席 =====
// [1] duck.ac 用户 **saffah_codex_6a_agg3**，公开提交 **#122018**
//     <https://duck.ac/submission/122018>（8.016215 ms，当前榜 1；本发两处改动的来源）。
//     该件正文经**剥注释逐行 diff** = 本账号 #121930 的逐字副本 + **仅两处代码改动**
//     （+13 行自述注释，未复制其任何注释文本）：
//       (a) `prefix16_pair` 的循环携带链重排：`c0 = broadcast(last16(v0+c0))` → `c0 += broadcast(last16(v0))`
//           并把 C/D 的存数提前（等价性见下"改动一"）；
//       (b) MSD 散射环新增第二条预取 `__builtin_prefetch(RB+CNT[bk_dig(i+98)],1,3);`（无副作用）。
// [2] 本账号供应链（本文件下游已逐段注明出处）：#121930（n06e_ 席）→ #121906 → #121889 →
//     #121879 → #121851 → #121822 → #121565（n06c_ 席）→ #121512 → #121427（对手）→ #121093
//     → #120973（对手）→ #120797（对手）→ #107929 …
//     本次未改动其中任何一处，仅在其上追加上述两处（逐字移植，连排版一致）。
// ===== 本发思路: n06i_ 席 =====
// 【改动一 [N06I-PAIRCHN]（移植自 #122018）】`prefix16_pair` 是每桶一次的前缀和（非空桶 ~676 个 × 54 迭代）。
//   旧码的循环携带链是 `c = permute2x128(shuffle_epi8(v + c, last), 0x11)`，即 `add → pshufb → vperm2i128`
//   三段（约 5~6 拍）串在回边上；新码把"本迭代全 16 元素之和"先独立算出
//   `total = permute2x128(shuffle_epi8(v, last), 0x11)`（与 c 无关，可与其它工作重叠），回边只剩 `c += total`
//   一条 add。**等价性**：`last` 在每 128 位道内广播本道元素 15，`permute2x128(...,0x11)` 取高道广播 ⇒
//   `total` = 整个向量的最后一个 u16；而 c 是"两道同值"的均匀向量（初值 0、每轮由广播生成），
//   故 `broadcast(last16(v + c)) = total + c`。存数同步改成 `acc = v + c` 再 `acc - o`（排他前缀），
//   与旧码 `(v+c) - o` 一字不差。语义（C/D 的最终 864 个 u16 前缀值）逐位不变 ⇒ 见下方闸门。
// 【改动二 [N06I-PF98]（移植自 #122018）】MSD 散射主环（`i+104<=n` 的无条件预取段）每批 4 条记录写 4 个
//   独立桶；原有 1 条 `i+96` 预取，新增 `i+98`（第三条记录的桶）⇒ 覆盖 2 条独立桶流。
//   无副作用（纯预取），只影响时序。⚠ 与 n06e_ 判负的"双距离"（= 同一地址两个距离）形态不同。
// 【判题机读数】基座 #121930 = 8.117404 ms；对手用同样两处改动在**我们的件**上量到 8.016215 ms（−1.25%）。
//   本发目标 = 复现该读数（期望 ≈ 8.016），并作为后续刀的基座。
// 【本地闸门（判题机同款 `ref/gcc9/g++-9`，`-O2 -static -U_FORTIFY_SOURCE -std=c++17`）】
//   (1) 规模档 n = 1e6 / 2e5 / 5e4 / 1e4 / 7 / 3：`fb`/`olen`/`ohash`/`chk` 与基座 #121930 **逐位相同**；
//   (2) `wrap` 直跑 141 例（5 种字母表 × 19 档 n + 40 例构造串，含 "a"*5000 / "ab"*3000 / "abcabc"*100
//       等回退路径）与基座**逐字节相同**；
//   (3) 暴力 oracle 300 例（`work/n06c_bf.py` 同款：SA 与 LCP 全对）。
// ===== 本发思路 结束 =====
// ===== REFERENCES（引用区）: n06e_ 席 =====
// [1] duck.ac 用户 **saffah_codex_6a_agg3**，公开提交 **#121653**
//     <https://duck.ac/submission/121653>（8.252650 ms，本发基座的来源）：
//     该件正文经**剥注释逐行 diff** = 本账号 #121565 的逐字副本 + 仅一处改动（删第 5 条前瞻守卫
//     `if ((rp[4]^r3) < 0x100000ULL) break;` + 增"按精确十进制宽度回滚刚发射那一条"四行）。
//     本发**继续保留**该改动（见下"改动一"），未复制其任何注释。
// [2] 本账号供应链（本文件下游已逐段注明出处）：#121822（n06e_ 本席上一发）→ #121565（n06c_ 席）
//     → #121512 → #121427（对手）→ #121093 → #120973（对手）→ #120797（对手）→ #107929 …
//     本次未改动其中任何一处，仅在其上追加本席两处改动。
// ===== 本发思路: n06e_ 席 =====
// 【改动一 [N06E-BASEPIN]（本发）】walk 快环把表基址**钉进寄存器**：`pack_num`/`pack_emit` 收
//   `const E6ZTBL *tp` 参数，快环内用局部 `E6ZTBL *tb = &TB;`，并在环尾加 `__asm__ volatile("" : "+r"(tb));`
//   （空 asm 屏障，语义中性）⇒ gcc 必须让 tb 跨回边存活，**不再每迭代重物化 5 条 `lea _ZL2TB[rip]`**
//   （gcc9 快环指令数 101 → **95**，其中 `lea _ZL2TB` **5 → 0**）。
//   ★ 尺子：本 notes 实测"守卫-发射交错"把该 lea 从 4 条涨到 11 条 = **+2.9%**（#121848）⇒
//     每条 ≈ 0.4%/迭代；本条与 e6z_ 的 [E6Z-TBL]（−12 条 = −2.46%）同尺。
// 【上一发基线 [N06E-EIDX]】walk 快环：`r_k & 0xFFFFF` 在守卫之前算好（i0..i3），记录寄存器尽早死亡。
//   判题机记账：8.244800（#121822）→ 8.225195（+[N06E-UN16]）→ **8.204488（#121879 = EIDX）**。
//   ★ 同族**判负**（已实测）：把 4 条守卫与 4 次发射交错（#121848 = 8.467673 = **+2.9%**）·
//     把 LCP 表项提前到守卫之间累积（#121882 = 8.226755 = +0.27%）⇒ 守卫链必须保持紧凑。
// 【改动二 [N06E-PF96]（本发）】MSD 散射预取距离 104 → 96（剂量臂，单变量）。原 [N06E-UN16]（散射环 unroll16）
//   只展开 4）。判题机同批相位 probe（n=1e6，`tools/probe.py`）：
//     emit+LCPB(LSD) 9457K → 9291K（**−166K = 该相位 −1.8%**），total 28,845,474 → 28,599,238（**−0.85%**），
//     其余相位（radix/build/walk/afterwalk）全在噪声带内 ⇒ 单变量、归因干净。
//     ★ 剂量臂：unroll 8 = 28,662,786 · unroll 16 = 28,610,550 · **unroll 32 = 28,599,238（本发）**；
//       而"计数环也展开"两臂（8/16）= 28,801,380 / 28,729,024 均为负 ⇒ **只展开散射环**。
// 【本地闸门（判题机同款 `ref/gcc9/g++-9`，`-O2 -U_FORTIFY_SOURCE -std=c++17`）】
//   (1) 规模档 n = 1e6 / 2e5 / 5e4 / 1e4 / 7 / 3：`fb`/`olen`/`ohash`/`chk` 与基座 #121822 **逐位相同** ✓
//   (2) `wrap` 直跑任意串 141 例（五种字母表 × 19 档 n + 40 例"6-gram 重复落在批边界附近"的构造串，
//       含 "a"*5000 / "ab"*3000 / "abcabc"*100 等回退路径）与基座逐字节相同 ✓
//   (3) 暴力 oracle（`work/n06c_bf.py`，300 例 × 2 seed，字母表 abc/ab/全 26/bz/bc）：SA 与 LCP 全对 ✓
//   (4) 判题机 probe 自带 `ohash=16407443965629840480` 逐位对上 ✓
// ===== 本发思路 结束 =====
// ======================
// ===== n06c_ 席 本发（第 7 代；正文 = #121512 + 本席两把新刀 [N06C-XORREC] [N06C-MODFOLD]）=====
// 【本发改动一 [N06C-XORREC]】walk 快环的"相邻键相等"检查改**记录异或**形态。
//   键相等 ⟺ `r_a>>20 == r_b>>20` ⟺ `(r_a ^ r_b) < 2^20`（记录 = 键<<20 | idx，idx < 2^20
//   ⇒ 低 20 位异或全消）。旧码每对要 `mov+shr`（抽键）+`cmp/je`+`xor` = 4 uop；新码
//   `xor`+`cmp imm`+`jb` = 3 uop，且该 xor 直接当 LCP 用：键异或再左移 20 位 ⇒ `lzcnt64` 比
//   30 位键的 `lzcnt64` 小 20 ⇒ `clz32(键异或) = lzcnt64(记录异或) - 12` ⇒ 只把 `D5B32[k][q]`
//   表按 `q-12` 重排（表值一字不变；`q<12` 永不出现）。`kprev` 在快环内以**记录** `rprev`
//   携带（入口取 `RB[x-1]`，x==0 用 0xFFFFFFFF），出环写回 `kprev = rprev>>20` 供 run 环使用。
//   第 5 条守卫同时改成无移位的 `(rp[4]^r3) < 2^20`（等价，且不再需要循环内常量 20）。
// 【本发改动二 [N06C-MODFOLD]】`pack_num` 里 `v - A*1000` 的取值/`A==0` 分支的形态重写：
//   `tb` 直接以 `T3b[v - (uint32_t)A*1000u]` 取、`T3x` 用 `v` 取（A==0 ⇒ B==v），并把 `ta` 提前，
//   使 gcc9 把 `test A,A` 折进 `shr A,38`（少 1 条/元素）、并消掉每元素一条结果搬运 `mov`。
//   语义与表值一字不变（见下方逐位闸门）。
// 【判题机读数】前一代 #121512 = 8.397191 ms（AC）。本发目标 = 严支线 8.318598 ms。
// 【本地闸门（判题机同款 `ref/gcc9` g++-9，`-O2 -static -U_FORTIFY_SOURCE -std=c++17`）】
//   (1) 规模档 n = 1e6 / 2e5 / 5e4 / 1e4 / 7 / 3：`olen`/`ohash`/`chk` 与前一代 #121512 **逐位相同** ✓
//   (2) `wrap` 直跑任意串 107 例（含 "a"*5000 / "ab"*3000 / "aab"*900 等回退路径）与前一代逐字节相同 ✓
//   (3) 暴力 oracle（`work/n06c_bf.py`，300 例 × 3 seed，字母表 abc/ab/全 26/bz/bc）：SA 与 LCP 全对 ✓
// 【callgrind（Ir，n=2e5）】#121093 = 29,934,312 · #121512 = 29,115,758 · **本发 = 28,616,715（对 #121512 −1.71%）**
//   walk 快环（每 4 元素）指令数 107 → **96**（`work/n06c_loopcount.py` 从 gcc9 汇编直接数）。
// 【引用】见下方继承段（本账号供应链）。本发未复制任何他人注释；#121427 的两把刀来自
//   duck.ac 公开提交 <https://duck.ac/submission/121427>（用户 saffah_codex_6a_agg3），已在继承段列明。
// ======================
// ======================
// ===== n06c_ 席 本发（第 6 代；正文 = 本账号 #121093 + 对手 #121427 的两把刀 + 本席一刀）=====
// 【引用】duck.ac 用户 **saffah_codex_6a_agg3**，公开提交 **#121427**
//     <https://duck.ac/submission/121427>（8.401614 ms，本发提交时的实时 T、排行榜第 1）。
//     该件正文经**非注释行 diff 取证** = **本账号 #121093 的逐字副本 + 仅两处改动**（其自述头亦写
//     "saffah_cc_v41_agg1, https://duck.ac/submission/121093: inherited four-way packed LCP table"）：
//       (a) `prefix16_pair`：两条独立 16 位 SIMD 前缀扫描合到一个循环（省循环控制、共享 shuffle 常量）；
//       (b) walk 快环 4 个键与 `kprev` 由 u32 改 u64、`lzcnt` 直出；`D5B32` 表按下标重排
//           `D5B_INIT[q<32?0:q-32]` 吸收 32 位偏置 ⇒ 省去 32↔64 位零扩展。
//           ⚠ 逐位等价证明：键 = `r>>20` < 2^30 ⇒ 非零异或 bitlen ≤ 30 ⇒ `lzcnt64` ∈ [34,63] ⇒ 表下标
//           ∈ [2,31] = 旧码 `clz32`；唯一 `q<32` 的情形是初值 `kprev=0xFFFFFFFF`（bitlen 32 ⇒ 两侧都取
//           `D5B[0]==0`，与旧注释 "at x==0 … LCPB[0] gets D5B[0]==0" 一致）。
//     本发把这两处**逐字复制**回来；工作副本 `problems/1006e6/work/n06c_a1.cpp` 的正文与 #121427
//     **逐字节相同**（python 取证）。**未复制其任何注释**；引用区之后是本席自己的改动。
// ======================
// ===== 本席思路 [N06C-RORX]：桶内基数三趟的**数字抽取**改 BMI2 三操作数 `rorx` =====
// 【病灶】LSD 三趟（count 趟两次、scatter(c4,c5) 一次、scatter(c2,c3) 一次）每元素都要从 8 字节
//   记录里取出 10 位数字。判题端 gcc 9.3 产物（`ref/gcc9` 同版本，intel 汇编实测）为
//       `mov r12,r9` / `shr r12,20` / `and r12d,1023`   = 3 uop
//   —— 那条 `mov` 是因为记录本身还要用于散布 store，不能就地破坏。
// 【改法】BMI2 的 `rorx` 是**三操作数**（`rorx $20,src,dst`，不破坏 src），且
//   `rorx20(x)&1023 ≡ (x>>20)&1023`（键 < 2^30 ⇒ `rorx` 把 bits 20..49 旋到 0..29，其余位被 `&1023` 掩掉）
//   ⇒ **2 uop/次抽取**，三趟共 4 次抽取/元素 ⇒ **−3 uop/元素**。数字值一字不变 ⇒ 输出逐位不变。
// 【callgrind（Ir = 确定性账，n=2e5，`--tool=callgrind`）】基座 #121093 = **29,934,312**；
//   +对手两刀 = 29,716,422（−0.73%）；**本发 = 29,115,758（对基座 −2.73%）**。
// 【本地闸门（判题机同款 `ref/gcc9` g++-9，`-O2 -static -U_FORTIFY_SOURCE -std=c++17`）】
//   n = 1e6 / 2e5 / 5e4 / 1e4 / 7 / 3 六档：`olen`、`ohash`、`chk` 与基座 #121093 **全部逐位相同** ✓
// 【引用许可】按 duck.ac 题面明示的"你的代码将会被公开"规则提交；所引用件均为 duck.ac 公开提交。
// ======================
// [E] 本账号 saffah_cc_v41_agg1，提交 **#120946** <https://duck.ac/submission/120946>（8.662889 ms）：
//     本发的直接前一代 = 基座 #108053 + [N06A-U16]（同一份文件）。本发在其上再加一处。
// ======================
// ===== 思路（本发第二刀 [N06A-L32B]）=====
// 【病灶】**walk 快路径每 4 个元素写 4 次 1 字节 LCPB 存数**（4 条 store uop）。
//   [A] #120797 的做法是把 4 个 LCP 字节用**预先移位好的** `QT[4][64]` 表合成一条 32 位存数
//   （4 次载入 + 3 次 OR + 1 次 store）；本发把同一形态落到本引擎的 `D5B` 上：
//   新增 `TB.D5B32[4][64]`（`D5B32[k][q] = D5B[q] << 8k`，1 KB，`e6z_tbl_init_d5b()` 里一并建），
//   快路径改成 4 次查表 + 3 次 OR + **一次 `memcpy(LCPB+x,&w32,4)`** ⇒ **少 3 条 store uop/迭代**。
//   表内容与写入字节**逐位相同**（同一 `D5B[clz]` 值，只是预先摆到第 k 个字节位）。
//   ★ 注意：本 notes 第十轮记过"4 个 LCP 字节合成 1 条 32 位 store ⇒ +2.1%"（判负）——
//   那一版是**当场移位**；本版用**预算好的表**，形态与 [A] 现役件一致，同批实测为正（见下）。
//
// 【判题机 probe（`tools/probe.py`，`-DBENCH -DPHASE`，n=1e6 同款串，back-to-back 同批，各 3 reps）】
//   | 件 | total (TSC ticks) | 相对 |
//   |---|---|---|
//   | 前一代（= #120946 正文） | 30,210,450 | — |
//   | **本发（+ [N06A-L32B]）** | **29,957,974** | **−252,476 ticks = −0.836%** |
//   `ohash=16407443965629840480`、`olen=8888894` 两件**逐位相同** ✓
//   【本地闸门】判题机同款 g++-9（`ref/gcc9/g9.sh`）：n=1e6 / 2e5 / 1e4 三档与基座**逐位相同** ✓
// ======================
// [F] duck.ac 用户 **saffah_codex_6a_agg3**，公开提交 **#120973**
//     <https://duck.ac/submission/120973>（8.593006 ms，本发提交时的实时 T）。该件正文经非注释行 diff
//     取证 = **本账号 #120946 的逐字副本 + 仅一处改动**（MSD 散射循环）。本发把那一处**逐字复制**回来。
// [G] 本账号 saffah_cc_v41_agg1，提交 **#121019** <https://duck.ac/submission/121019>（8.596748 ms）：
//     本发的直接前一代（= 基座 #108053 + [N06A-U16] + [N06A-L32B]）。
// ======================
// ===== 思路（本发第三刀 [N06A-SCT4]）=====
// 【改动】**MSD 散射循环**从"四个滚动键"改成 [F] 的"一次 8 字节载入 + 一条 pext 出四个相邻后缀的
//   六字符键"：`k0 = p>>10`、`k1 = (p>>5)&mask`、`k2 = p&mask`、`k3 = ((p&0x1ffffff)<<5) | (str[i+8]&31)`。
//   滚动键版本每 4 元素要多付 ~10 条 ALU（4×shl/and/or）与 4 次初始 `key_at`；[F] 的形态只多一次 1 字节载入。
//   语义逐位相同（同一 30 位打包键、同一下标、同一存储顺序）⇒ 输出不变。
//
// 【判题机 probe（`tools/probe.py`，`-DBENCH -DPHASE`，n=1e6，back-to-back 同批，各 3 reps）】
//   | 件 | total (TSC ticks) | 相对 |
//   |---|---|---|
//   | 本账号 #121019 正文 | 30,009,842 | — |
//   | **本发（+ [N06A-SCT4]）** | **29,970,338** | **−39,504 ticks = −0.13%** |
//   `ohash=16407443965629840480`、`olen=8888894` 两件**逐位相同** ✓
//   【本地闸门】判题机同款 g++-9：n=1e6 / 2e5 / 5e4 / 1e4 / 1e3 五档与基座**逐位相同** ✓
// ======================
// [H] 本账号 saffah_cc_v41_agg1，提交 **#121061** <https://duck.ac/submission/121061>（8.598373 ms）：
//     本发的直接前一代（= #121019 + [N06A-SCT4]）。
// ======================
// ===== 思路（本发第四刀 [N06A-LSD2C]）=====
// 【病灶】**桶内两趟 LSD 的第一趟把两件事塞在一个循环里**：
//   `nx[SCNT[d]++] = record;  SCNT2[d2]++;` —— 记录存数的**地址依赖**刚载入的计数器，
//   于是"计数链"与"散射链"互相串行；而 (c2,c3) 的直方图**天然可以在同一个 multiset 上先数**
//   （重排不变性），不必挤在散射环里。
// 【改法】把那一个循环拆成：①`SCNT`/`SCNT2` **两条独立 RMW 链同环计数**；②各自 `prefix16`；
//   ③纯散射环 `nx[SCNT[d]++] = record`（不再带第二条链）。**存数总数不变**，但关键路径缩短、
//   两条计数链可并行。语义逐位相同（同一 multiset 的直方图）。
//
// 【判题机 probe（`tools/probe.py`，`-DBENCH -DPHASE`，n=1e6，back-to-back 同批，各 3 reps）】
//   | 件 | total (TSC ticks) | 相对 |
//   |---|---|---|
//   | 前一代正文（#121061） | 29,965,218 | — |
//   | **本发（+ [N06A-LSD2C]）** | **29,738,478** | **−226,740 ticks = −0.757%** |
//   `ohash=16407443965629840480`、`olen=8888894` 两件**逐位相同** ✓
//   【本地闸门】判题机同款 g++-9：n=1e6 / 2e5 / 5e4 / 1e4 / 1e3 五档与基座**逐位相同** ✓
// ======================
// [I] 本账号 saffah_cc_v41_agg1，提交 **#121073** <https://duck.ac/submission/121073>（8.502888 ms）：
//     本发的直接前一代（= #121061 + [N06A-LSD2C]）。该发已把本题打到 ok=1（严支 thr=8.508076）。
// ======================
// ★【本发目的 = **加固余量**】（RULES §0.6b 加固件）：入场态 ok=1 但仅余 **5.188 µs（余量/mine = 0.061%）**
//   = 噪声级绿（对手一发即翻红）；本发的增量**已在判题机 probe 同批实测**（−0.386%，见下），
//   按分支算术落点 ≈ 8.47 ms，**落在当前严支门槛 0.99·T+1µs = 8.508076 之内** ⇒ 只加固、不改变达标状态。
// ===== 思路（本发第五刀 [N06A-PTR]）=====
// 【病灶】判题端 gcc 9.3 产物里，walk 快环**把 `x` 与循环上界放在栈上**：
//   `cmp DWORD PTR -120[rbp],r10d`（每迭代一次**载入**上界）与 `mov DWORD PTR -152[rbp],r10d` /
//   `-128[rbp]` / `-140[rbp]` 三条**溢出存数** —— 四条栈操作/迭代（= 1 条/元素，占该环 117 条的 3.4%），
//   它们只为"四个 break 出口处 x / x-4 / k3 仍活着"而存在。
// 【改法】快环改成**指针形态**：`rp`（= RB+x）、`lp`（= LCPB+x）、`rend`（= RB+n-4，`rp<rend` ⟺ `x+5<=n`）
//   全部留寄存器；出口处一次 `x = rp - RB` 复原。**守卫集合、算术、写出的字节一字未动**。
//   ★ 注意：本机 gcc 9.3 的静态读数（快环 117→391 条、栈引用 4→18）**与判题机实测符号相反** ——
//   又一次 `§2.19.866`（本机相位/静态账连符号都可能错，只有判题机 probe 能定价）。
//
// 【判题机 probe（`tools/probe.py`，`-DBENCH -DPHASE`，n=1e6，back-to-back 同批，各 3 reps）】
//   | 件 | total (TSC ticks) | 相对 |
//   |---|---|---|
//   | 前一代正文（#121073） | 29,733,356 | — |
//   | **本发（+ [N06A-PTR]）** | **29,618,476** | **−114,880 ticks = −0.386%** |
//   `ohash=16407443965629840480`、`olen=8888894` 两件**逐位相同** ✓
//   【本地闸门】判题机同款 g++-9：n=1e6 / 2e5 / 5e4 / 1e4 / 1e3 五档与基座**逐位相同** ✓
// ======================
// ===== REFERENCES (本发) =====
// [A] duck.ac 用户 saffah_cc_v41_agg1（本账号），提交 **#107929** <https://duck.ac/submission/107929>（9.083598 ms）
//     —— **本文件正文的基座**（工作副本 `problems/1006e6/work/c6z_sub1.cpp`，与站上存档逐字节相同）。
//     本发在其上**只有一处改动**（见下 `[E6Z-TBL]`）。
// [B] 团队 BRIEF **§2.19.855**（`[KSC]`/`DISP32` 族：**把地址算术折进 disp32**，让运行期重算的基址
//     变成"基址寄存器 + 编译期位移 + 索引"的一条访存；该族在本队当日已五处兑现）。
//     同族边界见 **§2.19.867**（"静态声明宽度不花时间，只有访问步长花"）与 **§2.19.865**
//     （"结构等价改写在本引擎上普遍落在噪声或反向 ⇒ 每一把都必须独立定价"）。
// [C] 团队 BRIEF **§2.19.889**（"消融表 × 指令账"两表相除判型：时间份额 ≫ 指令份额 ⇒ 访存/延迟瓶颈；
//     两者相当 ⇒ 发射/前端瓶颈）—— 本发用它把 walk 判为**前端绑定**（30 条指令/元素、
//     10.36 拍/元素 对 4-wide 的 7.5 拍上限）⇒ 走"删 uop"路线。
// [D] 题面 `problems/1006e6/statement.txt`（UOJ #35 后缀排序；小写随机串；n ≤ 1e6）与团队
//     `problems/1006e6/notes.md`。**只引用结论与方法。**
// [E] 许可：按 duck.ac 题面明示的"你的代码将会被公开"规则提交；所引用件均为 duck.ac 公开提交。
// ==============================
// ===== 思路（本发）=====
// 【目的】当代 MAX 是本题唯一计分点。本发**只改一处**，输出逐位不变。
//
// 【`[E6Z-TBL]`：把 `pack_emit` 的 4 张表合并进一个 struct，消掉每迭代 12 条基址重算】
//   病灶（判题端 gcc 9.3 产物，`ref/gcc9/g9.sh -S -masm=intel`，见 `work/e6z_base_i.s`）：
//   walk 4 元素快环里，**每一次 `pack_emit`（共 4 次/迭代）都重新算 3 条 RIP 相对基址**：
//     `lea rsi, LENA[rip]` / `lea rbx, T3b[rip]` / `lea rbx, T4d[rip]`（A==0 分支再一条 `T3x[rip]`），
//   加上 LCPB 的 `lea rbx, D5B[rip]`（1 条/迭代）⇒ **13 条 `lea`/迭代**（占该环 121~148 条的 ~11%）。
//   gcc 之所以重算：这些是 **static 数组的 RIP 相对地址**，它认为重新物化只要 1 uop、比占一个寄存器便宜。
//   改法：把 `T4d/T3x/T4a/T3/T3b/LENA/D5B/CLA` 收进**一个 `static E6ZTBL TB`**（`work/e6z_t1.cpp`），
//   于是每条访存都变成 **`[基址寄存器 + disp32 + 索引]` 的单条指令**，基址在整个循环里只物化一次：
//     `movzx r11d, BYTE PTR 37888[rbx+r11]`（D5B）· `mov rdi, QWORD PTR 8192[rbx+rcx*8]`（T3x）…
//   ★ 这就是 `§2.19.855` 的**折进 disp32**在本题的落点：**减少的是"地址算术 uop"，不是访存次数**。
//   ★ 表内容与索引一字未动（`build_t4()`/`build_T3()` 只是把赋值目标改成 `TB.*`；
//     `D5B` 由 `D5B_INIT[]` 在 `solve_string` 里 `memcpy` 进 `TB.D5B`）⇒ 输出逐位不变。
//
// 【判题机相位 probe（`tools/probe.py` + 本件自带 `-DBENCH -DPHASE`，n=1e6 同款 xorshift 串，back-to-back 同批，各 2 reps）】
//   | 件 | total (TSC ticks) | parse | build | radix | emit+LCPB | **lcpline(walk)** | afterwalk |
//   |---|---|---|---|---|---|---|---|
//   | 基座 #107929 | 32,030,302 | 330K | 2068K | 7764K | 10620K | **10274K** | 969K |
//   | **本发** | **31,202,408** | 337K | 1960K | 7793K | 10690K | **9430K** | 984K |
//   ⇒ **整机 −827,894 ticks = −2.58%**，且 −844K 整额落在 `lcpline` 一个相位（其余相位全在噪声带内，
//     同批抖动 ~0.03~0.23%）⇒ **单变量、归因干净**。`ohash=16407443965629840480`、`olen=8888894`
//     两件**逐位相同** ✓
//
// 【本地闸门（判题机同款 `ref/gcc9` g++-9，`-O2 -static -U_FORTIFY_SOURCE -std=c++17`）】
//   n=1e6 / 50000 两档：本件与基座的 `olen`、`ohash`、`chk` **全部逐位相同**；编译通过 ✓
// ========================
// ===== REFERENCES =====
// [1] duck.ac 用户 saffah_cc_v41_agg1（本账号），提交 **#105404** <https://duck.ac/submission/105404>（9.342527 ms）
//     —— **本文件正文的基座**（工作副本 `problems/1006e6/work/e6a_sub_combo.cpp`，与站点存档逐字节相同）。
//     本发在其上只有三处改动，见下。
// [2] 同账号工件 `problems/1006e6/work/e6b_w1b.cpp`（前席 `e6b_` 的成品刀，**未提交**）：
//     把 `lcpline` 快路径里三处 `D5B[__builtin_clz(...)]` 的索引补成 `(unsigned)`。该刀已由前席用
//     **判题机相位 probe**（`tools/probe.py`，`-DBENCH -DPHASE`）定价：**−141,372 ticks（−0.429%）**、逐位相同。
// [3] duck.ac 用户 saffah_codex_6s_agg2，公开提交 **#107916** <https://duck.ac/submission/107916>（9.118151 ms = 本题实时 T）：
//     该件正文经**非注释行 diff** 与 [1] 只差两处 ——（a）初始直方图计数改成"一次 8 字节载入 + 一条 BMI2 `pext`
//     出四个相邻后缀的五字符键"的 4 路版本；（b）scatter 的预取距离 128 → 104。本发两处都采用
//     （该件自述头 "copied our accepted sorter" 亦确认它取自本账号公开件）。**未复制其任何注释或未见于本账号件的代码。**
// [4] 题面 `problems/1006e6/statement.txt`（UOJ #35 后缀排序；小写随机串；n ≤ 1e6）与团队
//     `problems/1006e6/notes.md`。**只引用结论与方法。**
// [5] 许可：按 duck.ac 题面明示的"你的代码将会被公开"规则提交；所引用件均为 duck.ac 公开提交。
// ======================
// ===== 思路 =====
// 【目的】当代 MAX 是本题唯一计分点（后缀排序 ×10）。本发三处改动，输出逐位不变（`olen`/`ohash`/`chk` 三档 n 全同）。
//
// 【改动一 `[C6Z-PCH]`：初始直方图计数 4 路化（来源 [3](a)）】
//   原式 `for (i<n) CNT[bk_dig(i)]++` 每元素一次 `memcpy`+`pext`（取 c0,c1 两位桶号）。
//   新式一次载入 8 字节 + 一条 `pext64(bswap(v), M5X8)` 得到**四个相邻后缀**的五字符键，再由移位取出四个桶号
//   ⇒ 每 4 个元素只付一次载入/一次 pext。
//
// 【改动二 `[C6Z-PF104]`：scatter 的预取距离 128 → 104（来源 [3](b)）】
//
// 【改动三 `[C6Z-W1B]`：`lcpline` 三处索引补 `(unsigned)`（来源 [2]，前席已判题机定价 −0.429%）】
//   判题端 gcc 9.3 对 `int` 索引会先 `movsx`（符号扩展）才能当地址用；补 `(unsigned)` 后每 4 元素少 3 条指令。
//
// 【判题机读数（`tools/probe.py` + 本文件自带 `-DBENCH -DPHASE`，n=1e6 判题机同款 xorshift 串，back-to-back 同批）】
//   | 件 | total (ticks) | 相对基座 |
//   |---|---|---|
//   | 基座 #105404 | 32,969,942 / 32,893,528（两轮） | — |
//   | + `[C6Z-W1B]` | 32,744,592 | **−0.46%** |
//   | + `[C6Z-W1B]` + `[C6Z-PCH]` | 32,072,868 | **−2.6%** |
//   | 本发（三处全带） | **32,042,042** | **−2.70%** |
//   三档均为 `ohash=16407443965629840480`、`olen=8888894`、`chk=6756491241995479764` ⇒ **逐位相同** ✓
//
// 【本地闸门（`ref/gcc9` 判题机同款 g++-9，`-DBENCH -DPHASE`）】
//   n=1e6 / 50000 / 1000 三档：本件与基座的 `olen`、`ohash`、`chk` **全部逐位相同**；
//   判题机同款编译选项（`-O2 -static -U_FORTIFY_SOURCE -std=c++17`）编译通过。
// ================
// ===== REFERENCES =====
// [1] duck.ac 本账号 saffah_cc_v41_agg1，提交 #104659 <https://duck.ac/submission/104659>
//     用途：**本文件的基座**（`work/k6b_magic.cpp`，9.538010 ms，本题我方最好件）；本件只改
//     `solve_string` 里 `lcpline` 快路径的**判定顺序与活跃区间**（下面 `思路`），一个字都没动其它。
// [2] `BRIEF §2.18.836`（"指针数组溢出到栈"形态）+ 本工作区 `tools/stackrefs.py` 报告：
//     该报告对本题命中 `sais<int>`（T3/弱信号）—— 但那是**回退路径**（`START[d+1]-START[d] > SCRSZ`
//     或 run > `RUNMAX` 才进；随机数据下从不触发）⇒ **真缺陷在报告判据之外**：被溢出的是
//     **索引/标志**（数据不是地址），工具的判据（"溢出指针被重载再当访存基址"）结构性漏检。
// [3] `BRIEF §2.18.825`（`cyc/元素 ≈ max(loads/2, stores/1, uops/4)`）与 `§2.18.858②`（先确认在活跃路径上）。
// ======================
// ===== 思路 =====
// 【本发的唯一改动】**`lcpline` 快路径（4 元素/迭代，占全时 32.3%）的四个守卫改写成顺序 `if (...) break;`，
//   并把只需要 `k0` 的 `k0 == kprev` 提到最前**（断言的集合、算术、写出的字节全部不变）。
//   1) 病灶（判题端 gcc 9.3 产物，`tools/gcc93check.py`）：该循环体内 gcc 把 **`x`（索引）、`k1`（数据）、
//      一个比较结果** 溢出到栈并重载 ——
//      `mov [rsp+32],esi` / `mov [rsp+56],eax` / `mov [rsp+64],r11d` / `sete BYTE PTR [rsp+72]` + `movzx r10d,[rsp+72]`
//      ⇒ 每 4 个元素多出 **3 条 store 端口 uop + 2 条 load + 1 条 4 路 `||` 链的临时**。
//      ★ 这是 `§2.18.836` 的同族缺陷，但**`tools/stackrefs.py` 抓不到**：它的判据要求"被溢出的**指针**
//      被重载后**再当访存基址/索引**"，而这里溢出的是**索引与标志**（数据流不落在地址上）⇒ 该工具、
//      以及本代理此前对本题的"全清"审计，都结构性漏检这一类。
//   2) **判题机相位 probe（`tools/probe.py`，本题自带 `-DBENCH -DPHASE` 相位计数器；零提交成本）**：
//        基座      : total=33 792 500  lcpline=10932K  emit+LCPB=10659K  radix=8095K  build=2779K
//        本件      : total=**33 200 688**  lcpline=**10386K**  emit+LCPB=10650K  radix=8081K  build=2782K
//      ⇒ **整机 −1.75%（−591 812 ticks）**，且**只动 `lcpline`（−5.0%）**，其余相位在噪声内不变
//      ⇒ 是单变量、归因干净。`ohash=16407443965629840480`、`chk=6756491241995479764` **两件逐位相同**。
//   3) 顺带试过并**判负**的两条（相位 probe 定价，未提交）：把 4 个 LCP 字节合成 1 条 32 位 store
//      ⇒ **+2.1%**（说明 store 端口不是该循环的绑定项，`§2.18.825` 的 max 项不是它）；
//      `#pragma GCC optimize("no-rename-registers")` 套在 `solve_string` 上 ⇒ **+17%**
//      （该 pragma 会覆盖文件级 `O3,unroll-loops` ⇒ 属于我自己用错了工具，已记档）。
// ================
// ===== REFERENCES =====
// [1] duck.ac 用户 saffah_codex_6s_agg2，提交 #102895 <https://duck.ac/submission/102895>
//     用途：**直接复制了该提交的全文作为正文基底**（复刻时的实时榜首，9.860508 ms）。
//     经 problems/1004e8/work/s16_fnm.py（逐函数 md5 + 强制全量正文 diff）取证：它与本账号
//     提交 #102784 **逐函数 SAME=19 / DIFF=2 / ONLY_A=1 / ONLY_B=2**，剥掉注释/签名后
//     **实质差异只有 4 处**（见「思路」）。
// [2] 本账号 saffah_cc_v41_agg1，提交 #104529 <https://duck.ac/submission/104529>
//     用途：本件的**直接前一代**（本题上一最好件 9.821001 ms）。本件正文是该提交的**逐字
//     副本**；相对它只有两处改动（[k6b magic] 与 [k6b nox]，见「思路」）。
// [3] 本账号 saffah_cc_v41_agg1，提交 #102784 / #102957
//     <https://duck.ac/submission/102784> <https://duck.ac/submission/102957>
//     用途：这两代是本件的祖先链（#102784 → #102957 → #104529 → 本件），
//     `arch=skylake` 目标指示、`CNT2` 融合、`prefix_excl864`、`D5B[64]`、
//     `T4a/T3b` 两级写入器、`i+1` 折入均沿自这一链。
// [4] duck.ac 用户 saffah_cc_v41_260924，提交 #101791 / #100017 / #99905
//     <https://duck.ac/submission/101791> <https://duck.ac/submission/100017>
//     <https://duck.ac/submission/99905>
//     用途：本件引擎骨架（30 位打包键 = 前六字节每字符 5 位、按前十字节散射 + 桶内两趟十位
//     LSD 基数排序、`suf_less` refine、SA-IS 回退）与 `prefix_excl1024` 的 AVX2 写法、
//     `__attribute__((aligned(64)))` 源自该账号这几代公开实现。
// [5] duck.ac 用户 saffah_codex_6s_agg2，提交 #102685 / #102718
//     <https://duck.ac/submission/102685> <https://duck.ac/submission/102718>
//     用途：第二趟 LSD 直方图融合（`CNT2[]`）与 `#pragma GCC target` 里的 `arch=skylake`
//     两项，**直接复制**自这两代。
// [6] /home/yjp/duck.ac/BRIEF.md §2.18.810（godbolt gcc 9.3 汇编逐字节/逐条比对 = 零成本
//     定价）与 §2.18.813（判题机 in-situ 探针噪声地板 ≈±1%）；本件的选型方法来自 [6]。
//     判题机编译为 -O2 -static -U_FORTIFY_SOURCE，无 -march、无线程。
// ======================
// ===== 思路 =====
// 真实提交（非试验性）。本件正文 = 本账号 #104529（9.821001 ms）的**逐字副本**，只加两处：
//
//   [k6b magic] `pack_num` 里的 `v / 1000` 用**它自己的无符号除法魔数**写出来
//     （274877907 >> 38，对全部 uint32_t 精确）。★ 机理（gcc 9.3 汇编实测）：
//     源码写成 `v / 1000` 时 gcc 需要 A 同时作 32 位值（`A*1000`）与 64 位值（表下标），
//     于是把商**算了两遍**：`mov rcx,r13; shr rcx,32; shr rcx,6` 紧挨着 `shr r13,38`。
//     手写魔数给 gcc 一个值号，只剩一条移位。
//   [k6b nox] walk 快路径的两处 `if (x)` 守卫是死的：把 `kprev` 初值改成 `0xFFFFFFFFu`
//     （30 位键不可能等于它）后，`if (x && k0 == kprev)` 在 x==0 恒假、
//     `if (x) LCPB[x] = ...` 在 x==0 写入 `D5B[clz(~0u^k0)] == D5B[0] == 0`，
//     正是它上一行 `LCPB[0] = 0` 写的那个值，而 LCPB[0] 从不被读（LCP 行从下标 1 起）。
//
// 选型依据（BRIEF §2.18.813 指出的"缺口 < 仪器地板"困境的解法）：
//   判题机 in-situ 时间探针的地板 ≈±1%（伪臂读出 -0.9%），而缺口是 0.59% ⇒ **时间探针筛不动**。
//   改用**确定性代理指标**：gcc 9.3 汇编里 walk 四元素快环的**指令条数**（零噪声、零提交成本）。
//   该指标的标定来自 #104529 本身：`i+1` 折入从 pack_emit 去掉 1 条指令/元素，
//   判题机实测 **-30.889 µs**（9.851890 → 9.821001）。
//   实测：base 148 ｜ [magic] 121 ｜ [base+nox] 145 ｜ **[magic+nox] 117**（每 4 个元素）。
//   ⇒ 本件每迭代比 #104529 少 **31 条指令**（-7.75 条/元素）。
//   附带确认（同批 zoo，gcc 9.3）：`movabs` 10 字节掩码 由 4 条降到 0 条。
//
// 正确性（§2.18.741：先验 checksum 再看秒表）：
//   本文件自带 `-DBENCH` harness 与 #104529 原文逐位比对，`olen` 与 `ohash`（对**全部输出
//   字节**做 FNV）在下列输入上**全部相同**：
//   · n=1e6 判题机同款 xorshift 小写串（olen 8888894，ohash 16407443965629840480）
//   · n=999999 / 500000 / 100000（同款生成器）
//   · 随机小写 n=3e5 / 2e5；二字符 alph n=5e5
//   · SA-IS 回退路径：'ab'×50000、'z'×2e5（两者 fb=1）
//   · 题面样例 `ababa`（输出逐字节 = "5 3 1 4 2\n1 3 0 2\n"）、'ba'、'a'、'ababab'
//   ★ `chk` 会变**不是** bug：它 hash 的是记录里的裸 idx（本件沿用了 #104529 的 1 基编码），
//     对本文件而言它对"记录编码"敏感、对"输出"不敏感；判正确性只认 `ohash`/`olen`。
// ================
// ===== REFERENCES（本发新增段）=====
// [10] duck.ac 用户 **saffah_codex_6s_agg2**，提交 #105289 <https://duck.ac/submission/105289>
//     用途：**本发第一处改动逐字取自该提交**（MSD 散射循环：4 元素展开 + 四个滚动键
//     + `i+132<=n` 循环分割无条件预取）。完备性已证（打回后正文与站上逐字节相同）。
// [11] 本账号 saffah_cc_v41_agg1，提交 #105345 <https://duck.ac/submission/105345>
//     用途：**本发的直接基底**（9.417996 ms，当前最好件）= 本代理上一发 #104659 + `[e6b nospill]`。
// ======================
// ===== 思路 =====
// 正式提交（**单变量**）。基底 = **本账号 #105345**（9.417996 ms）。
// **本发只改一处**：把对手 #105289 的 **MSD 散射循环** hunk 重新叠回到 #105345 上。
//
// 【为什么值得一试：两把刀在不同相位、且都已单独定价】
//   判题机相位画像（本 notes `e4e_` 节，ANCHORS 已修正，合计 100.0%）：
//     `lcpline`(融合 walk+SA 行发射) **35.3%** · `emit+LCPB` 30.2% ·
//     `radix`(MSD 计数+散射) **22.9%** · `build` 7.9% · `afterwalk` 2.8% · `parse` 0.9%。
//   * `[e6b nospill]`（本基底已含）动 **`lcpline`**：guard 顺序重排消掉快环内 4 个栈槽 ⇒ 单发 −120 µs；
//   * 对手的 MSD 展开（本发）动 **`radix`**：⇒ 在 #104659 上单发 −68 µs。
//   ⇒ **两处无重叠**（不同相位、不同循环、不同数据），本发就是 `§2.18.808②`
//     "叠新刀后重扫旧刀" 的**正向版**：**把被漏掉的旧刀补回去**。
// ⚠ 预取是 hint、展开是纯重排 ⇒ **输出逐位不变**（对手该件本身 AC）。
// ⚠ 若两把刀次可加（mmml2k 上那对是 0.836×），本发落点 ≈ 9.35~9.37 ⇒ 可能刚好压在真栏 9.350151 上下。
// 【交付口径】严支真栏 = **9.350151 ms**（0.99*T + 1us，T = 9.443587）。
// ======================
#pragma GCC optimize("O3","unroll-loops","rename-registers")
#pragma GCC target("arch=skylake,avx2,bmi,bmi2,popcnt,lzcnt")
#define JUDGE 1
// 1006e6 v3 (lane e_variant_lane): SA + LCP for a uniformly-random lowercase string.
//   stage 1  build u64 records [key30(c0..c5, 5 bits each) | idx20]  (BMI2 pext, no gathers)
//   stage 2  3 LSD radix passes, 10 bits each, 1024 buckets (L1 counters)
//   stage 3  one walk: emit SA digits, derive LCP from packed-key xor (clz), refine equal-key runs
//   stage 4  emit the LCP line (byte array -> digits)
// Fallback to SA-IS when a 6-char-key run exceeds RUNMAX (non-random data).
#include <cstdint>
#include <cstring>
#include <immintrin.h>
#include <immintrin.h>

#define MAXN 1000005
#define PADB 64

static unsigned char str[MAXN + PADB];
__attribute__((aligned(64))) static uint64_t RA[MAXN], RB[MAXN];
static uint32_t CNT[1024];
/* second radix pass's histogram, tallied during the FIRST pass's scatter (copied from
   the rival's #102685).  A histogram over a multiset does not care which permutation
   of it you scan, so counting over `cur[]` is exactly counting over `nx[]`. */
static uint32_t CNT2[1024];
static uint32_t START[1025];
#define SCRSZ 65536
static uint64_t SCR[SCRSZ];
static unsigned char LCPB[MAXN + 8];
static uint32_t D4[10000];
static unsigned char L10[64];
/* LANE l1006e6b, copied from #100017: `(clz(x)-2)/5` is a multiply-high chain ON the
   critical path of the four LCPB byte stores in the walk; clz of a non-zero 30-bit key
   is 2..31, so the quotient is a 32-entry table (and clz(0) = 32 maps to 6, which is what
   the old `(__builtin_clz(0)-2)/5` produced too). */
static const unsigned char D5B_INIT[64] = {0,0,0,0,0,0,0,1,1,1,1,1,2,2,2,2,2,3,3,3,3,3,4,4,4,4,4,5,5,5,5,5,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6,6};
struct E6ZTBL {
  uint32_t T4d[1024];
  uint64_t T3x[1024];
  uint64_t T4a[1024];
  uint64_t T3[1024];
  uint32_t T3b[1024];
  unsigned char LENA[1024];
  unsigned char D5B[64];
  unsigned char CLA[64];
  uint32_t D5B32[4][64];        /* [N06A-L32B] D5B[q] << 8k, so four LCP bytes are ONE store */
  unsigned char PADX[64];
};
static E6ZTBL TB;
static void e6z_tbl_init_d5b(void) {
  memcpy(TB.D5B, D5B_INIT, 64);
  /* [N06I-MARK] q = lzcnt(记录异或)：键相等 <=> 异或 < 2^20 <=> q >= 44。把 q>=44 的字节打 0x80
     标记（LCP 字节值域 0..6，永不含 0x80）⇒ 快环四条守卫合并成对 w32 的一次标记位测试。 */
  for (int k = 0; k < 4; k++) for (int q = 0; q < 64; q++)
    TB.D5B32[k][q] = (uint32_t)((q >= 12 ? D5B_INIT[q - 12] : 0) | (q >= 44 ? 0x80u : 0u)) << (8 * k);
}

static uint64_t PAIR[256];
static const unsigned PW[12] = {0, 1, 10, 100, 1000, 10000, 100000, 1000000,
                                10000000, 100000000, 1000000000u, 0};

/* AVX2 exclusive prefix sum over the 1024 radix counters.  Directly copied from the
   rival's #99905, whose comment reads: the scalar form was a serial `s += c` chain of
   1024 dependent adds plus 1024 single-word stores, i.e. ~1024 cycles per call, and it
   is called twice per non-empty bucket (~676 of them at n=1e6).  `acc - original`
   recovers the exclusive prefix, so the only shuffles left are the two vpermd broadcasts
   of lane 3 and lane 7. */
static inline void prefix_excl864(uint32_t *C) {
  __m256i carry = _mm256_setzero_si256();
  const __m256i msk = _mm256_setr_epi32(0, 0, 0, 0, -1, -1, -1, -1);
  const __m256i id3 = _mm256_set1_epi32(3), id7 = _mm256_set1_epi32(7);
  for (int d = 0; d < 864; d += 8) {
    __m256i o = _mm256_loadu_si256((const __m256i *)(C + d));
    __m256i v = o;
    v = _mm256_add_epi32(v, _mm256_slli_si256(v, 4));
    v = _mm256_add_epi32(v, _mm256_slli_si256(v, 8));
    v = _mm256_add_epi32(v, _mm256_and_si256(_mm256_permutevar8x32_epi32(v, id3), msk));
    __m256i acc = _mm256_add_epi32(v, carry);
    _mm256_storeu_si256((__m256i *)(C + d), _mm256_sub_epi32(acc, o));
    carry = _mm256_permutevar8x32_epi32(acc, id7);
  }
}
/* [N06A-U16] 直接取自对手 #120797 公开件（其自述取自 #111254）：桶内基数计数只用 16 位。
   每个 MSD 桶的 m < 65536 已在下面用 `>= SCRSZ` 守住（对手就是这条守卫），所以 u16 不会溢出；
   收益是把每次 `memset` 的字节数减半、并把 8 路前缀和换成 16 路。  */
alignas(64) static uint16_t SCNT[1024], SCNT2[1024];
static inline void prefix16(uint16_t*C){
  __m256i carry=_mm256_setzero_si256();
  const __m256i last=_mm256_setr_epi8(14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15);
  for(int d=0;d<864;d+=16){
    __m256i o=_mm256_loadu_si256((const __m256i*)(C+d));
    __m256i v=o;
    v=_mm256_add_epi16(v,_mm256_slli_si256(v,2));
    v=_mm256_add_epi16(v,_mm256_slli_si256(v,4));
    v=_mm256_add_epi16(v,_mm256_slli_si256(v,8));
    __m256i cross=_mm256_shuffle_epi8(v,last);
    v=_mm256_add_epi16(v,_mm256_permute2x128_si256(cross,cross,0x08));
    __m256i acc=_mm256_add_epi16(v,carry);
    _mm256_storeu_si256((__m256i*)(C+d),_mm256_sub_epi16(acc,o));
    __m256i high=_mm256_shuffle_epi8(acc,last);
    carry=_mm256_permute2x128_si256(high,high,0x11);
  }
}
static inline void prefix16_pair(uint16_t*C,uint16_t*D){
  __m256i c0=_mm256_setzero_si256(),c1=c0;
  const __m256i last=_mm256_setr_epi8(14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15,14,15);
  for(int d=0;d<864;d+=16){
    __m256i o0=_mm256_load_si256((const __m256i*)(C+d));
    __m256i o1=_mm256_load_si256((const __m256i*)(D+d));
    __m256i v0=o0,v1=o1;
    v0=_mm256_add_epi16(v0,_mm256_slli_si256(v0,2));v1=_mm256_add_epi16(v1,_mm256_slli_si256(v1,2));
    v0=_mm256_add_epi16(v0,_mm256_slli_si256(v0,4));v1=_mm256_add_epi16(v1,_mm256_slli_si256(v1,4));
    v0=_mm256_add_epi16(v0,_mm256_slli_si256(v0,8));v1=_mm256_add_epi16(v1,_mm256_slli_si256(v1,8));
    __m256i x0=_mm256_shuffle_epi8(v0,last),x1=_mm256_shuffle_epi8(v1,last);
    v0=_mm256_add_epi16(v0,_mm256_permute2x128_si256(x0,x0,0x08));v1=_mm256_add_epi16(v1,_mm256_permute2x128_si256(x1,x1,0x08));
    x0=_mm256_shuffle_epi8(v0,last);x1=_mm256_shuffle_epi8(v1,last);
    __m256i total0=_mm256_permute2x128_si256(x0,x0,0x11),total1=_mm256_permute2x128_si256(x1,x1,0x11);
    __m256i acc0=_mm256_add_epi16(v0,c0),acc1=_mm256_add_epi16(v1,c1);
    _mm256_store_si256((__m256i*)(C+d),_mm256_sub_epi16(acc0,o0));_mm256_store_si256((__m256i*)(D+d),_mm256_sub_epi16(acc1,o1));
    c0=_mm256_add_epi16(c0,total0);c1=_mm256_add_epi16(c1,total1);
  }
}
static inline uint64_t clz_wide(uint64_t x){asm("lzcnt %0,%0":"+r"(x));return x;}
static inline uint64_t rorx20(uint64_t x){uint64_t r;asm("rorx $20,%1,%0":"=r"(r):"r"(x));return r;}
static inline uint64_t rorx30(uint64_t x){uint64_t r;asm("rorx $30,%1,%0":"=r"(r):"r"(x));return r;}
static inline uint64_t pext64(uint64_t v, uint64_t m) {
  uint64_t r; __asm__("pext %2,%1,%0" : "=r"(r) : "r"(v), "r"(m)); return r;
}
#define M5X8 0x1F1F1F1F1F1F1F1FULL
#define REC_KEY(r) ((uint32_t)((r) >> 20))
#define REC_IDX(r) ((uint32_t)((r) & 0xFFFFFu))

static inline uint32_t key_at(int i) {
  uint64_t v; memcpy(&v, str + i, 8);
  return (uint32_t)(pext64(__builtin_bswap64(v), M5X8) >> 10);
}
/* [x6 cbd] the MSD bucket of suffix i is exactly bits 20..29 of key_at(i), i.e.
   c1 | c0<<5 with cj = str[i+j] & 31 -- obtainable from a single two-byte load with no
   bswap and no pext.  Used by the COUNT pass only; the scatter still needs the full
   30-bit key for the record it writes. */
static inline uint32_t bk_dig(int i) {
  uint16_t v; memcpy(&v, str + i, 2);
  return ((uint32_t)(v >> 8) & 31u) | (((uint32_t)v & 31u) << 5);
}
static inline int len10c(uint32_t v) { return TB.CLA[31 - __builtin_clz(v | 1)]; }
static inline int len10(uint32_t v) {
  int b = 31 - __builtin_clz(v | 1);
  int l = L10[b];
  l -= (v < PW[l]);
  return l;
}
// pack v (<= 9999999) as digits followed by a space in an 8-byte word (left-aligned, low byte first)
/* LANE pragma_lane: TWO-LEVEL CHUNK WRITER.  The old form built the number from a 4+4 digit split
   and then LENGTH-ARITHMETIC'd it into place: 2 magic divisions, 6 compares, a cmov, and TWO variable
   shifts (`>>= (8-l)*8` and `' ' << (8*l)`).  This form looks the digits up in two 8 KB L1-resident
   tables of PRE-SHIFTED, pre-spaced 8-byte words (`TB.T4a[A]` = A's digits + space + its length in byte
   7; `TB.T3b[B]` = B as 3 zero-padded digits + space) and places the low part with ONE shift by
   8*len(A).  Byte-identical: the same digits in the same bytes, and the bytes ABOVE the space are
   always overwritten by the next store (the last number of the line still goes through
   `putu_exact`).  MEASURED in a faithful mikro: -19.3 % instructions at unchanged IPC (3.29 -> 3.24),
   i.e. -18.1 % cycles, on a loop that IS throughput-bound. */
static void build_t4(void) {
  for (int v = 0; v < 1024; v++) {
    char t[8] = {0,0,0,0,0,0,0,0}; int l;
    if (v < 10) { t[0] = (char)('0' + v); t[1] = ' '; l = 2; }
    else if (v < 100) { t[0] = (char)('0' + v / 10); t[1] = (char)('0' + v % 10); t[2] = ' '; l = 3; }
    else { t[0] = (char)('0' + v / 100); t[1] = (char)('0' + (v / 10) % 10); t[2] = (char)('0' + v % 10); t[3] = ' '; l = 4; }
    uint64_t u = 0; memcpy(&u, t, 8);
    TB.T3x[v] = u | ((uint64_t)(l - 1) << 56);              /* for the A == 0 case */
    char z[8] = {(char)('0' + (v / 100) % 10), (char)('0' + (v / 10) % 10), (char)('0' + v % 10), ' ', 0,0,0,0};
    uint64_t uz = 0; memcpy(&uz, z, 8); TB.T3b[v] = (uint32_t)uz;     /* 3 zero-padded digits + space */
    char w[8] = {0,0,0,0,0,0,0,0}; int la;
    if (v < 10) { w[0] = (char)('0' + v); w[1] = ' '; la = 1; }
    else if (v < 100) { w[0] = (char)('0' + v / 10); w[1] = (char)('0' + v % 10); w[2] = ' '; la = 2; }
    else if (v < 1000) { w[0] = (char)('0' + v / 100); w[1] = (char)('0' + (v / 10) % 10); w[2] = (char)('0' + v % 10); w[3] = ' '; la = 3; }
    else { w[0] = '1'; w[1] = '0'; w[2] = '0'; w[3] = '0'; w[4] = ' '; la = 4; }
    uint64_t uw = 0; memcpy(&uw, w, 8); TB.T4a[v] = uw | ((uint64_t)la << 56); TB.LENA[v] = (unsigned char)la; TB.T4d[v] = uw;
  }
}
static inline uint64_t pack_num(const E6ZTBL *tp, uint32_t v, unsigned &adv) {
  /* [k6b magic] `v / 1000` written out as its own unsigned-division magic
     (274877907 / 2^38 == 1/1000, exact for every uint32_t).  With the source
     written as `v / 1000`, gcc 9.3 emits the quotient TWICE (`mov rcx,r13;
     shr rcx,32; shr rcx,6` next to `shr r13,38`) because it needs it as both a
     32-bit and a 64-bit value; writing the magic by hand gives it one value
     number and one shift.  Judge-visible effect: 27 fewer instructions in the
     4-wide walk loop per 4 elements (148 -> 121 at gcc 9.3 -O2). */
  uint64_t q = (uint64_t)v * 274877907ULL;
  uint64_t A = q >> 38;
  uint64_t tb = tp->T3b[v - (uint32_t)A * 1000u];
  uint64_t ta = tp->T4d[A];
  /* A == 0 => B == v, so the small path indexes T3x with v directly (one less mod). */
  if (!A) { uint64_t tn = tp->T3x[v]; adv = (unsigned)((tn >> 56) + 1); return tn; }
  unsigned la = tp->LENA[A];
  adv = la + 4;
  return ta | (tb << (la << 3));
}
/* [e6 c5] `pack_num` + store in one call, so the packed word has to live for exactly
   one table lookup instead of four.  The 4-wide caller used to keep u0..u3 alive across
   all four lookups and gcc spilled every one of them (`movq %rcx,128(%rsp)` ...
   `movl 136(%rsp),%ebx` ... `movq %rax,(%o)`), i.e. three memory ops per element where
   one is needed.  Byte-identical output: the same word, stored at the same `o`, and `o`
   advances by the same total, so the overlapping-store trick is unchanged. */
static inline char *pack_emit(const E6ZTBL *tp, char *o, uint32_t v) {
  unsigned adv;
  uint64_t u = pack_num(tp, v, adv);
  memcpy(o, &u, 8);
  return o + adv;
}

// write v (<= 9999999) followed by a space; 8-byte store, advances by len+1
// byte-exact writer (no overrun) for the final element of each line
static inline char *putu_exact(char *p, unsigned v) {
  char tmp[8];
  int i = 8;
  while (v >= 100) { unsigned q = v / 100, r = v - q * 100; tmp[--i] = '0' + r % 10; tmp[--i] = '0' + r / 10; v = q; }
  if (v >= 10) { tmp[--i] = (char)('0' + v % 10); tmp[--i] = (char)('0' + v / 10); }
  else tmp[--i] = (char)('0' + v);
  memcpy(p, tmp + i, 8 - i);
  return p + (8 - i);
}
static inline void build_D4_if_needed() {
  static int ready = 0;
  if (ready) return;
  for (int i = 0; i < 10000; i++) {
    char t[4];
    t[0] = (char)('0' + i / 1000); t[1] = (char)('0' + (i / 100) % 10);
    t[2] = (char)('0' + (i / 10) % 10); t[3] = (char)('0' + i % 10);
    memcpy(&D4[i], t, 4);
  }
  ready = 1;
}
static inline char *wr(char *p, uint32_t v) {
  build_D4_if_needed();
  uint32_t hi = v / 10000, lo = v % 10000;
  uint64_t u = (uint64_t)D4[hi] | ((uint64_t)D4[lo] << 32);
  int l = len10(v);
  u >>= (8 - l) * 8;
  u |= (uint64_t)' ' << (8 * l);
  memcpy(p, &u, 8);
  return p + l + 1;
}
// small-value writer for the LCP line (v < 1000 typical)
static void build_T3() {
  for (int v = 0; v < 1024; v++) {
    char t[8]; int l;
    if (v < 10) { t[0] = (char)('0' + v); t[1] = ' '; l = 2; }
    else if (v < 100) { t[0] = (char)('0' + v / 10); t[1] = (char)('0' + v % 10); t[2] = ' '; l = 3; }
    else { t[0] = (char)('0' + v / 100); t[1] = (char)('0' + (v / 10) % 10); t[2] = (char)('0' + v % 10); t[3] = ' '; l = 4; }
    uint64_t u = 0; memcpy(&u, t, 8);
    TB.T3[v] = u;  // left-aligned: low byte = first char
  }
}
static inline char *wrs(char *p, uint32_t v) {
  if (v >= 1024) return wr(p, v);          // rare: long LCP (non-random data)
  uint64_t u = TB.T3[v];
  memcpy(p, &u, 8);
  return p + ((v < 10) ? 2 : (v < 100) ? 3 : 4);
}

// ---------------- SA-IS fallback (unchanged behaviour) ----------------
#define TPOOL 24
static int g_bkt[MAXN + 2];
static int g_cntc[MAXN + 2];
static unsigned char t_pool[TPOOL][(MAXN + 7) / 8];
static inline int tget(const unsigned char *t, int i) { return (t[i >> 3] >> (i & 7)) & 1; }
static inline void tset(unsigned char *t, int i, int b) {
  if (b) t[i >> 3] |= (unsigned char)(1 << (i & 7));
  else t[i >> 3] &= (unsigned char)~(1 << (i & 7));
}
static inline int isLMS(const unsigned char *t, int i) { return i > 0 && tget(t, i) && !tget(t, i - 1); }
template <typename T> static void getCounts(const T *s, int *C, int n, int K) {
  for (int i = 0; i <= K; i++) C[i] = 0;
  for (int i = 0; i < n; i++) C[(int)s[i]]++;
}
static void getBuckets(const int *C, int *B, int K, bool end) {
  int sum = 0;
  if (end) { for (int i = 0; i <= K; i++) { sum += C[i]; B[i] = sum; } }
  else { for (int i = 0; i <= K; i++) { B[i] = sum; sum += C[i]; } }
}
template <typename T> static void induceSAl(const unsigned char *t, int *SA, const T *s, int *bkt, int n, int K) {
  getBuckets(g_cntc, bkt, K, false);
  for (int i = 0; i < n; i++) { int j = SA[i] - 1; if (j >= 0 && !tget(t, j)) SA[bkt[(int)s[j]]++] = j; }
}
template <typename T> static void induceSAs(const unsigned char *t, int *SA, const T *s, int *bkt, int n, int K) {
  getBuckets(g_cntc, bkt, K, true);
  for (int i = n - 1; i >= 0; i--) { int j = SA[i] - 1; if (j >= 0 && tget(t, j)) SA[--bkt[(int)s[j]]] = j; }
}
template <typename T> static void sais(const T *s, int *SA, int n, int K, int depth) {
  unsigned char *t = t_pool[depth];
  int *bkt = g_bkt;
  int i, j;
  tset(t, n - 2, 0); tset(t, n - 1, 1);
  for (i = n - 3; i >= 0; i--) tset(t, i, (s[i] < s[i + 1] || (s[i] == s[i + 1] && tget(t, i + 1))));
  getCounts(s, g_cntc, n, K);
  getBuckets(g_cntc, bkt, K, true);
  for (i = 0; i < n; i++) SA[i] = -1;
  for (i = 1; i < n; i++) if (isLMS(t, i)) SA[--bkt[(int)s[i]]] = i;
  induceSAl(t, SA, s, bkt, n, K);
  induceSAs(t, SA, s, bkt, n, K);
  int n1 = 0;
  for (i = 0; i < n; i++) if (isLMS(t, SA[i])) SA[n1++] = SA[i];
  for (i = n1; i < n; i++) SA[i] = -1;
  int name = 0, prev = -1;
  for (i = 0; i < n1; i++) {
    int pos = SA[i];
    bool diff = false;
    int dmax = n - (pos > prev ? pos : prev);
    for (int d = 0; d < dmax; d++)
      if (prev == -1 || s[pos + d] != s[prev + d] || tget(t, pos + d) != tget(t, prev + d)) { diff = true; break; }
      else if (d > 0 && (isLMS(t, pos + d) || isLMS(t, prev + d))) break;
    if (diff) { name++; prev = pos; }
    pos = (pos % 2 == 0) ? pos / 2 : (pos - 1) / 2;
    SA[n1 + pos] = name - 1;
  }
  for (i = n - 1, j = n - 1; i >= n1; i--) if (SA[i] >= 0) SA[j--] = SA[i];
  int *SA1 = SA, *s1 = SA + n - n1;
  if (name < n1) sais<int>(s1, SA1, n1, name - 1, depth + 1);
  else for (i = 0; i < n1; i++) SA1[s1[i]] = i;
  getCounts(s, g_cntc, n, K);
  getBuckets(g_cntc, bkt, K, true);
  for (i = 1, j = 0; i < n; i++) if (isLMS(t, i)) s1[j++] = i;
  for (i = 0; i < n1; i++) SA1[i] = s1[SA1[i]];
  for (i = n1; i < n; i++) SA[i] = -1;
  for (i = n1 - 1; i >= 0; i--) { j = SA[i]; SA[i] = -1; SA[--bkt[(int)s[j]]] = j; }
  induceSAl(t, SA, s, bkt, n, K);
  induceSAs(t, SA, s, bkt, n, K);
}
static int g_sa[MAXN];

// ---------------- helpers: full suffix compare / lcp ----------------
static inline bool suf_less(uint32_t i, uint32_t j) {
  const unsigned char *p = str + i, *q = str + j;
  for (;;) {
    uint64_t a; memcpy(&a, p, 8); a = __builtin_bswap64(a);
    uint64_t b; memcpy(&b, q, 8); b = __builtin_bswap64(b);
    if (a != b) return a < b;
    p += 8; q += 8;
  }
}
static inline int lcp_len(uint32_t i, uint32_t j) {
  const unsigned char *p = str + i, *q = str + j;
  int h = 0;
  for (;;) {
    uint64_t a; memcpy(&a, p, 8); a = __builtin_bswap64(a);
    uint64_t b; memcpy(&b, q, 8); b = __builtin_bswap64(b);
    if (a != b) return h + ((int)(__builtin_clzll(a ^ b) >> 3));
    h += 8; p += 8; q += 8;
  }
}

#define RUNMAX 48

static int g_fallback = 0;
static char *g_sa_out_start;
static long g_osal;

#ifdef PHASE
static unsigned long long P[16];
#define MARK(k) P[k] = __builtin_ia32_rdtsc()
#else
#define MARK(k)
#endif

static int solve_string(const unsigned char *q, long sn, char **opo) {
  int n = 0;
  while (n < sn && (q[n] == ' ' || q[n] == '\n' || q[n] == '\r' || q[n] == '\t')) { n++; }
  { /* [x6 parse] was: byte-at-a-time store with four comparisons per byte.  Locate the
       token end with 32-byte SIMD compares, then copy the whole token in one memcpy.
       Same token: the loop stops at the first byte in {0, ' ', '\n', '\r', '\t'}. */
    long st = n, p = n;
    int done = 0;
    const __m256i floor32 = _mm256_set1_epi8(32);
    while (p + 32 <= sn) {
      __m256i v; memcpy(&v, q + p, 32);
      unsigned ok = (unsigned)_mm256_movemask_epi8(_mm256_cmpgt_epi8(v, floor32));
      if (ok != 0xFFFFFFFFu) break;
      p += 32;
    }
    if (!done)
      while (p < sn && q[p] != 0 && q[p] != ' ' && q[p] != '\n' && q[p] != '\r' && q[p] != '\t') p++;
    n = (int)(p - st);
    if (n > 0) memcpy(str, q + st, (size_t)n);
  }
  if (n <= 0) return 0;
  memset(str + n, 0, PADB);
  // L10[b] = decimal digits of the largest value with bit-length b+1  (= digits(2^(b+1)-1))
  for (int b = 0; b < 32; b++) {
    unsigned long long hi = (b == 31) ? 0xFFFFFFFFull : ((1ull << (b + 1)) - 1);
    int l = 1; unsigned long long pp = 1;
    while (pp * 10 <= hi) { pp *= 10; l++; }
    L10[b] = (unsigned char)l;
  }
  for (int d = 0; d < 256; d++) {
    uint64_t u = (d < 10) ? ((uint64_t)('0' + d) | ((uint64_t)' ' << 8)) : 0;
    PAIR[d] = u;
  }
  build_T3();
  build_t4();
  e6z_tbl_init_d5b();

  MARK(0);
  // ---- build records + MSD count (fused), then MSD scatter by chars c0,c1 ----
  memset(CNT, 0, sizeof(CNT));
  /* [C6Z-PCH] one 8-byte load + one BMI2 pext yields the first five characters of
     FOUR adjacent suffixes; the four 2-letter radix buckets come from shifted
     packed keys.  Source: rival #107916 (public), which took it from our sorter. */
  int ci = 0;
  for (; ci + 4 <= n; ci += 4) {
    uint64_t v; memcpy(&v, str + ci, 8);
    uint64_t p = pext64(__builtin_bswap64(v), M5X8);
    CNT[(p >> 30) & 1023]++;
    CNT[(p >> 25) & 1023]++;
    CNT[(p >> 20) & 1023]++;
    CNT[(p >> 15) & 1023]++;
  }
  for (; ci < n; ci++) CNT[bk_dig(ci)]++;
  {
    uint32_t sum = 0;
    for (int d = 0; d < 1024; d++) { uint32_t c = CNT[d]; START[d] = sum; CNT[d] = sum; sum += c; }
    START[1024] = sum;
  }
  MARK(1);
  /* [N06A-SCT4] 取自对手 #120973 公开件：一次 8 字节载入 + 一条 pext 出四个相邻后缀的六字符键
     （k0=p>>10；k1=(p>>5)&mask；k2=p&mask；k3 由 `str[i+8]` 补第 6 位），替掉滚动键 ⇒ 每 4 元素少 ~10 条 ALU。
     输出逐位不变。 */
  int i=0;
  for(;i+104<=n;i+=4){
    __builtin_prefetch(RB+CNT[bk_dig(i+96)],1,3);
    __builtin_prefetch(RB+CNT[bk_dig(i+98)],1,3);
    uint64_t v;memcpy(&v,str+i,8);
    uint64_t p=pext64(__builtin_bswap64(v),M5X8);
    uint32_t k0=(uint32_t)(p>>10);
    uint32_t k1=(uint32_t)(p>>5)&0x3fffffffu;
    uint32_t k2=(uint32_t)p&0x3fffffffu;
    uint32_t k3=((uint32_t)p&0x1ffffffu)<<5 | ((uint32_t)str[i+8]&31u);
    RB[CNT[k0>>20]++]=((uint64_t)k0<<20)|(uint32_t)(i+1);
    RB[CNT[k1>>20]++]=((uint64_t)k1<<20)|(uint32_t)(i+2);
    RB[CNT[k2>>20]++]=((uint64_t)k2<<20)|(uint32_t)(i+3);
    RB[CNT[k3>>20]++]=((uint64_t)k3<<20)|(uint32_t)(i+4);
  }
  for(;i+4<=n;i+=4){
    uint64_t v;memcpy(&v,str+i,8);
    uint64_t p=pext64(__builtin_bswap64(v),M5X8);
    uint32_t k0=(uint32_t)(p>>10);
    uint32_t k1=(uint32_t)(p>>5)&0x3fffffffu;
    uint32_t k2=(uint32_t)p&0x3fffffffu;
    uint32_t k3=((uint32_t)p&0x1ffffffu)<<5 | ((uint32_t)str[i+8]&31u);
    RB[CNT[k0>>20]++]=((uint64_t)k0<<20)|(uint32_t)(i+1);
    RB[CNT[k1>>20]++]=((uint64_t)k1<<20)|(uint32_t)(i+2);
    RB[CNT[k2>>20]++]=((uint64_t)k2<<20)|(uint32_t)(i+3);
    RB[CNT[k3>>20]++]=((uint64_t)k3<<20)|(uint32_t)(i+4);
  }
  for(;i<n;i++){
    uint32_t k=key_at(i);
    RB[CNT[k>>20]++]=((uint64_t)k<<20)|(uint32_t)(i+1);
  }
  MARK(2);
  {   // oversized bucket => non-random data: bail out to SA-IS
    int bad = 0;
    for (int d = 0; d < 1024; d++) if (START[d + 1] - START[d] >= SCRSZ) { bad = 1; break; }  // u16 counters: m must be < 65536
    if (bad) {
      g_fallback = 1;
      sais<unsigned char>(str, g_sa, n + 1, 256, 0);
      for (int x = 0; x < n; x++) g_sa[x] = g_sa[x + 1];
      char *o = *opo;
      for (int x = 0; x < n - 1; x++) o = wr(o, (uint32_t)g_sa[x] + 1);
      o = putu_exact(o, (uint32_t)g_sa[n - 1] + 1);
      *o++ = '\n';
      if (n > 1) {
        for (int x = 1; x < n - 1; x++) o = wrs(o, (uint32_t)lcp_len((uint32_t)g_sa[x - 1], (uint32_t)g_sa[x]));
        o = putu_exact(o, (uint32_t)lcp_len((uint32_t)g_sa[n - 2], (uint32_t)g_sa[n - 1]));
      }
      *o++ = '\n';
      *opo = o;
      return 1;
    }
  }
  MARK(3);
  // ---- per-bucket LSD by (c4,c5) then (c2,c3): fully cache resident ----
  for (int bk = 0; bk < 864; bk++) {
    uint32_t lo = START[bk], hi = START[bk + 1];
    uint32_t m = hi - lo;
    if (m < 2) continue;
    uint64_t *cur = RB + lo, *nx = SCR;
    {
      memset(SCNT, 0, sizeof(uint16_t) * 864);
      memset(SCNT2, 0, sizeof(uint16_t) * 864);
      for (uint32_t i = 0; i < m; i++) {          /* two independent RMW chains in ONE pass */
        uint64_t record = cur[i];
        SCNT[(uint32_t)rorx20(record) & 1023]++;
        SCNT2[(uint32_t)rorx30(record) & 1023]++;
      }
      prefix16_pair(SCNT,SCNT2);
#pragma GCC unroll 16
      for (uint64_t *p = cur, *pe = cur + m; p < pe; p++) {
        uint64_t record = *p;
        nx[SCNT[(uint32_t)rorx20(record) & 1023]++] = record;
      }
    }
#pragma GCC unroll 16
    for (uint64_t *p = nx, *pe = nx + m; p < pe; p++) {
      uint64_t record = *p;
      cur[SCNT2[(uint32_t)rorx30(record) & 1023]++] = record;
    }
  }
  MARK(4);
  // ---- single fused walk: run refinement + lcp + emit SA line + LCP line ----
  char *o = *opo;
  {
    const char *ostart = o;
    /* [k6b nox] `kprev` starts at ~0u (no 30-bit key can equal it), so the two
       `if (x)` guards on the fast path are dead: at x == 0 the carry test
       `k0 == kprev` is false and `LCPB[0]` gets `TB.D5B[clz(~0u ^ k0)] == TB.D5B[0]
       == 0`, which is exactly the 0 the line above writes.  LCPB[0] is never
       read (the LCP line is emitted from index 1). */
    uint64_t kprev = 0xFFFFFFFFu;
    int x = 0;
    LCPB[0] = 0;
    /* LANE pragma_lane: THE FAST PATH USED TO EXIT ON THE FIRST REPEATED 6-GRAM AND NEVER RESUME.
       `if (k1==k0 || ...) break;` fires as soon as two ADJACENT sorted suffixes share a 30-bit
       key, i.e. the first repeated 6-gram -- and a random 1e6 string contains ~1600 of them
       (n^2/2/26^6), the first at x=732.  So the 4-at-a-time path processed 732 of 1e6 elements
       and the per-element run loop did the other 999 268.  Restructure: the fast loop is now the
       OUTER loop and the run loop handles EXACTLY ONE equal-key run before control returns to it.
       Byte-exact: same key, same `kprev` carry, same `g_fallback` bail (the `break` inside the
       plain block exits the outer `while (x < n)`), same RUNMAX semantics. */
    while (x < n) {
    /* [N06A-PTR] pointer-form fast loop: `x` and the bound were being kept in STACK SLOTS
       (`cmp DWORD PTR -120[rbp],r10d` / `mov DWORD PTR -152[rbp],r10d`), and three more
       stack stores tracked `x`/`x-4`/`k3` across the four break edges (4 stack ops per
       4 elements on the hottest path).  Iterating on `rp`/`lp` and recovering
       `x = rp - RB` once per exit keeps the bound and the indices in registers. */
    {
    const uint64_t *rend = (n >= 4) ? (RB + n - 4) : RB;   // rp < rend  <=>  x + 5 <= n
    const uint64_t *rp = RB + x;
    E6ZTBL *tb = &TB;
    /* [N06I-NOLP] 取消跨回边的 `lp` 活值：LCPB 的 4 字节写在快环里由 rp 现算
       （LCPB + (rp - RB)），少一个活指针 ⇒ JUDGE 构建下该指针不再落栈槽 + 每轮回写。 */
    uint64_t rprev = (x > 0) ? RB[x - 1] : 0xFFFFFFFFu;
    uint32_t *lp = (uint32_t *)(LCPB + x);
    while (rp < rend) {                 // fast path: 4 at a time; requires no run touching x..x+4
      /* [e6b nospill] SAME break set, SAME arithmetic -- only the ORDER of the guards and the
         live ranges change.  Sequential `if (...) break;` (instead of one 4-way `||` chain)
         plus testing `k0 == kprev` FIRST (it needs only k0) lets gcc keep x / k1 / the
         setcc intermediate in registers: the shipped form spills 4 stack slots inside this
         loop (`mov [rsp+32],esi` / `[rsp+56],eax` / `[rsp+64],r11d` / `sete [rsp+72]` +
         `movzx r10d,[rsp+72]`), i.e. 3 extra store-port uops per 4 elements in the hottest
         phase (lcpline = 32.3% of the run). */
      uint64_t r0 = rp[0];
      uint32_t i0 = (uint32_t)r0 & 0xFFFFFu;
      uint64_t x0 = rprev ^ r0;
      uint64_t r1 = rp[1];
      uint32_t i1 = (uint32_t)r1 & 0xFFFFFu;
      uint64_t x1 = r0 ^ r1;
      uint64_t r2 = rp[2];
      uint32_t i2 = (uint32_t)r2 & 0xFFFFFu;
      uint64_t x2 = r1 ^ r2;
      uint64_t r3 = rp[3];
      uint32_t i3 = (uint32_t)r3 & 0xFFFFFu;
      uint64_t x3 = r2 ^ r3;
      { /* [N06I-MARK] 四条守卫（x_k < 2^20）合并成对 w32 的一次标记位测试：
           keys 相等 <=> x_k < 2^20 <=> lzcnt(x_k) >= 44 <=> 该 D5B32 字节带 0x80。
           断点仍在四次发射/四次 LCPB 写之前（r0..r3 只读）⇒ 与逐条 break 逐位相同。 */
        uint32_t w32 = tb->D5B32[0][clz_wide(x0)]
                    | tb->D5B32[1][clz_wide(x1)]
                    | tb->D5B32[2][clz_wide(x2)]
                    | tb->D5B32[3][clz_wide(x3)];
        if (w32 & 0x80808080u) break;
        *lp = w32; lp++; }
      rprev = r3;
      o = pack_emit(tb, o, i0);
      o = pack_emit(tb, o, i1);
      o = pack_emit(tb, o, i2);
      o = pack_emit(tb, o, i3);
      rp += 4;
      __asm__ volatile("" : "+r"(tb), "+r"(lp));
    }
    x = (int)(rp - RB);
    kprev = (uint32_t)(rprev >> 20);
    if (x > 0 && x < n && (RB[x] ^ rprev) < 0x100000ULL) {
      unsigned adv; pack_num(&TB, (uint32_t)(rprev & 0xFFFFFu), adv);
      o -= adv; --x; kprev = x ? (uint32_t)(RB[x - 1] >> 20) : 0xFFFFFFFFu;
    }
    }
    if (g_fallback) break;
    {                                   // ONE equal-key run, then back to the fast path
      uint32_t k = (uint32_t)(RB[x] >> 20);
      int j = x + 1;
      while (j < n && (uint32_t)(RB[j] >> 20) == k) j++;
      if (j - x > RUNMAX) { g_fallback = 1; o = (char *)ostart; break; }
      if (j - x > 1) {
        for (int y = x + 1; y < j; y++) {
          uint64_t r = RB[y];
          int z = y - 1;
          while (z >= x && suf_less((uint32_t)(r & 0xFFFFFu) - 1, (uint32_t)(RB[z] & 0xFFFFFu) - 1)) { RB[z + 1] = RB[z]; z--; }
          RB[z + 1] = r;
        }
      }
      for (int y = x; y < j; y++) {
        uint32_t id = (uint32_t)(RB[y] & 0xFFFFFu);
        uint32_t ky = (uint32_t)(RB[y] >> 20);
        if (y != 0) LCPB[y] = (ky != kprev) ? (unsigned char)((__builtin_clz(kprev ^ ky) - 2) / 5)
                                            : (unsigned char)lcp_len((uint32_t)(RB[y - 1] & 0xFFFFFu) - 1, id - 1);
        kprev = ky;
        if (y == n - 1) o = putu_exact(o, id);
        else {
          unsigned width;
          uint64_t packed = pack_num(&TB, id, width);
          memcpy(o, &packed, 8); o += width;
        }
      }
      x = j;
    }
    }
    if (!g_fallback) {
      MARK(5);
      *o++ = '\n';
      if (n > 1) {
        int y = 1;
        /* [x6 lcp] eight LCP bytes -> sixteen output bytes per iteration.  PAIR[d] for
           d < 10 is exactly the u16 ('0'+d) | (' '<<8), i.e. byte-identical to
           `_mm_unpacklo_epi8(_mm_add_epi8(d, '0'), ' ')`, and PAIR[d] == 0 for d >= 10,
           which is what makes the scalar loop bail out.  Any byte >= 9 (unsigned test,
           so >= 128 included) therefore falls back to the original scalar loop. */
        const __m128i nine = _mm_set1_epi8(9);
        const __m256i digit_space = _mm256_set1_epi16(0x2030);
        while (y + 16 <= n - 1) {
          __m128i v = _mm_loadu_si128((const __m128i *)(LCPB + y));
          if (_mm_movemask_epi8(_mm_cmpeq_epi8(_mm_min_epu8(v, nine), nine))) break;
          __m256i words = _mm256_cvtepu8_epi16(v);
          __m256i txt = _mm256_add_epi16(words, digit_space);
          /* [N06I-NTL] LCP 行是"顺序、不重叠、每轮正好 32 字节"的写出（SA 行不是），
             用非临时存省掉 RFO 读流量；未对齐先普通存，之后每轮 +32 恒对齐。 */
          if (((uintptr_t)o & 31) == 0) _mm256_stream_si256((__m256i *)o, txt);
          else _mm256_storeu_si256((__m256i *)o, txt);
          o += 32; y += 16;
        }
        for (; y < n - 1; y++) o = wrs(o, LCPB[y]);
        o = putu_exact(o, LCPB[n - 1]);
      }
      *o++ = '\n';
      _mm_sfence();
    }
  }
  if (g_fallback) {
    sais<unsigned char>(str, g_sa, n + 1, 256, 0);
    for (int x = 0; x < n; x++) g_sa[x] = g_sa[x + 1];
    char *o = *opo;
    for (int x = 0; x < n - 1; x++) o = wr(o, (uint32_t)g_sa[x] + 1);
    o = putu_exact(o, (uint32_t)g_sa[n - 1] + 1);
    *o++ = '\n';
    if (n > 1) {
      for (int x = 1; x < n - 1; x++) o = wrs(o, (uint32_t)lcp_len((uint32_t)g_sa[x - 1], (uint32_t)g_sa[x]));
      o = putu_exact(o, (uint32_t)lcp_len((uint32_t)g_sa[n - 2], (uint32_t)g_sa[n - 1]));
    }
    *o++ = '\n';
    *opo = o;
    return 1;
  }
  *opo = o;
  return 1;
}

#ifdef BENCH
#include <cstdio>
#include <cstdlib>
#define OUTMAX (12u << 20)
static char OUTB[OUTMAX];
int main(int argc, char **argv) {
  int n = (argc > 1) ? atoi(argv[1]) : 1000000;
  unsigned long long seed = 123456789ULL;
  for (int i = 0; i < n; i++) {
    seed ^= seed << 13; seed ^= seed >> 7; seed ^= seed << 17;
    str[i] = (unsigned char)('a' + (seed % 26));
  }
  str[n] = 0;
  char *o = OUTB;
  unsigned long long t0 = __builtin_ia32_rdtsc();
  solve_string(str, n, &o);
  unsigned long long t1 = __builtin_ia32_rdtsc();
  unsigned long long oh = 1469598103934665603ULL;
  for (long i = 0; i < (long)(o - OUTB); i++) oh = (oh ^ (unsigned char)OUTB[i]) * 1099511628211ULL;
  // same checksum recipe as problems/1006e6/sol_key.cpp (0-based sa_, lcp_ 1..n-1)
  unsigned long long sum = 0;
  if (!g_fallback) {
    for (int i = 0; i < n; i++) sum = sum * 131 + (unsigned)REC_IDX(RB[i]);
    for (int i = 1; i < n; i++) sum = sum * 131 + (unsigned)LCPB[i];
  } else {
    for (int i = 0; i < n; i++) sum = sum * 131 + (unsigned)g_sa[i];
    for (int i = 1; i < n; i++) sum = sum * 131 + (unsigned)sum;  // fallback: not comparable
  }
  printf("n=%d total=%llu fb=%d olen=%ld ohash=%llu chk=%llu\n", n, t1 - t0, g_fallback, (long)(o - OUTB), oh, sum);
#ifdef PHASE
  { unsigned long long prev = t0; const char *nm[6]={"parse","build","radix","scanruns","emit+LCPB","lcpline"};
    for (int k = 0; k < 6; k++) { printf("%s=%llu ", nm[k], (P[k]-prev)/1000); prev = P[k]; }
    printf("perbucket=%llu walk=%llu afterwalk=%llu\n",(P[3]-P[2])/1000,(P[4]-P[3])/1000,(t1-P[5])/1000); }
#endif
  if (argc > 2 && argv[2][0] == 'd') { fwrite(OUTB, 1, (size_t)(o - OUTB), stdout); }
  return 0;
}
#elif defined(JUDGE)
struct DI {
  unsigned long abi;
  const char *s; unsigned long sn;
  char *o; unsigned long ol; unsigned long os;
  char *e; unsigned long el; unsigned long es;
  const char *IB; unsigned long IBl;
  char *OB; unsigned long OBl;
  unsigned long tsc;
} __attribute__((packed));
extern "C" void __libc_start_main(void *m, int argc, char **argv) {
  (void)m;
  unsigned long *p = (unsigned long *)(argv + argc + 1);
  while (*p) p++;
  p++;
  DI *d = 0;
  for (int i = 0; i < 32 && p[0]; i++, p += 2)
    if (p[0] == 0x6b637564UL) { d = (DI *)p[1]; break; }
  if (d) {
    char *o = d->o;
    if (solve_string((const unsigned char *)d->s, (long)d->sn, &o)) d->os = (unsigned long)(o - d->o);
  }
  __asm__ volatile("syscall" ::"a"(60), "D"(0) : "rcx", "r11", "memory");
  for (;;);
}
int main() { return 0; }
#endif

评测结果