// ===== 本次思路 s4（p02g_ 席：= #121708 的正文 + **单变量** DFS deg-2 链 MLP）=====
// 【动机】#121686（sort+链+soff 三族同批）= 26.776453 比 #121598 慢 1.6%，无法归因；#121708
// （只留 sort 两处 uop 削减）= **26.193065**（比 #121598 快 0.605%）⇒ sort 族已判正，链/soff 族待判。
// 本件在 #121708 正文上**只加链一处**（soff 不动，因它是 s2 中唯一会改变 DFS 寄存器分配的改动）。
// 【本刀】deg-2 链折叠：原式在 w 已知处只发 `prefetch(&PS(w))`，真正的 `sv = PS(v)` 载入在**循环体
// 末尾**才做 ⇒ 该随机 L3 载入与其后的 `adj[sv]` 载入串行。本刀把**真载入**提到 prefetch 处
// （cut-edge 判定与 emit(w) 都不碰 sv/sv1；break 时旧值原样保留），并从提前的值发 adj[] 预取。
// 判题机探针 #3（同窗、臂间机器码相同、空对照 0.096%）定价 **−0.15 ~ −0.24%**；探针 #6 差分 −0.267%。
// 【闸门】与 #121708 逐位等价：396 随机小图 fails 0 · 十形状×三尺度全 SAME · 官尺度 tc19(n=5e5) 与
// shape110(n=5e5) 输出 md5 逐字节同。
// ===== REFERENCES (p02g_ seat, 2026-10-02) =====
// ===== 本次思路 s3（p02g_ 席：只保留 sort 环的两处 uop 削减，用于**单变量板面判决**）=====
// 【动机】本席 s2（#121686 = 26.776453）把 sort 环 uop 削减、DFS deg-2 链 MLP、emit soff 局部化
// 三族**同批**发出 ⇒ 板面 +1.6% 无法归因。本件只留：
// (1) sort 环交换判据：`cmpq pair,swapped` 的 CF 恰为 `lo>hi`（相等时为 0，与 bsort len==2 特化同谓词），
// `movl d,tt; adcl $0,tt` ⇒ `tt==3` 严格等价 `d==2 && lo>hi`（本支只在 d<3 进入）⇒ 一条 `cmpl $3; cmove`；
// (2) 叶子队列：题面明示图**连通**且 m ≤ n ⇒ 本环的 m==n 情形是单圈连通图，度 ≥ 1 ⇒
// `d < 2` 与 `d == 1` 等价 ⇒ `cmpl $2; adcl $0,qt`。
// 两者合计 cachegrind 确定性 Ir 118 125 844 → 117 3xx xxx（≈ −0.7%），误预测不变。
// 【闸门】与 #121598 逐位等价：396 随机小图 fails 0 · 十形状×三尺度全 SAME · 官尺度 tc19(n=5e5)
// 与 shape110(n=5e5) 输出 md5 逐字节同。
// ===== REFERENCES (p02g_ seat, 2026-10-02) =====
// ===== REFERENCES (p02g_ seat, 2026-10-02) =====
// [G1] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121532** <https://duck.ac/submission/121532>
// （判题机 Time = 26.337815 ms = 本题当前 T，已 Accepted）—— 本文件从该件取用**两处技术**：
// (a) **分隔符查表压缩**（`pairread`）：把 `la=ctz(mask) / end=ctz(mask&(mask-1)) / lb=end-la-1`
// 的两次依赖 `ctz` + 下标算术，换成一次 **8 KiB `u16` 表** `dtab[mask & 8190]` 载入 ——
// 低 4 位存 `end`、高 12 位存 `(la*8+lb)*16`（即 `PSH` 的 16B 块偏移）。
// 表项按 `delim_code((x<<1)|8192)` 常量求值生成（`std::array` + `constexpr`），
// 「bit0 恒 0（`++p` 后首字节是数字）· bit13 恒 1（标签 ≤ 6 位 ⇒ 第二分隔符必在头 14 字节内）」
// 两条输入不变量把索引压到 13 位 ⇒ 表 4096×2B = 8 KiB（替 32 KiB 的 PSH 直索引）。
// (b) **短表排序的 cmov 交换**（`d < 3` 支）：把 `u64 pair = *(u64*)(adj+b)` 一次载入/写回，
// 字节旋转得 `swapped`，`change = (d==2) & (lo32 > hi32)`，再一条
// `testl/cmovnz` 选择 —— 替本文件原「min/max + 两条掩码合并」的四步形态。
// 其件正文 = 本账号 **#121280** 正文 + 恰好上述 **2 处**语义差（本席 `g++ -fpreprocessed -dD -E`
// 剥注释后逐行 diff 核证：预处理器输出仅 2 个 hunk，其余为注释）。原提交正文按站点规则公开可见，
// 无独立许可证声明；本引用仅取该两处技术，合规见 RULES.md 第 3 节。
// [G2] duck.ac 用户 **saffah_cc_v41_agg1**（本账号）提交 **#121280** <https://duck.ac/submission/121280>
// （26.586786 ms）—— 本文件基座（含本账号 p02f_ 席的排序派发无分支化），其自带的历史引用/思路
// 注释原样保留在下方。
// ===== 本次思路（p02g_ 席，抄件 + 单变量叠刀，逐位等价）=====
// 【入场】`mine = 26.586786 (#121280)` · `T = 26.337815 (#121532，新于我件 ⇒ 严支 0.99T+1µs = 26.075437，
// 缺 511.349 µs = 1.923%；宽支 1.005T = 26.470504 亦缺 116.282 µs)` ⇒ 翻红。
// 【本发 s1】只做**取回**：把 #121532 相对本账号 #121280 的 2 处语义差原样叠回（单变量对照 = 纯抄件）。
// 叠件等价性论证：(a) 两条输入不变量（见 [G1](a)）在本题数据格式下恒成立，且 `x=0`/单比特两种
// 退化掩码在表生成里显式返回 0；(b) `change` 谓词 `x > y` 与 `bsort` len==2 特化的 `x < y` 在
// 「相等」情形同为不交换 ⇒ 逐位同；d==1 时 `change==0` ⇒ 写回原值，与掩码形态同。
// ===== REFERENCES (p02f_ seat, 2026-10-02) =====
// [F1] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121117** <https://duck.ac/submission/121117>
// （27.185966 ms）—— peel 前瞻距离 `qq[qh + 8]` → **`qq[qh + 12]`**（本文件已含；其件与 #121037 的唯一差别）。
// [F2] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121037** <https://duck.ac/submission/121037>（27.203961 ms）
// —— 本文件正文的基座（= 本账号 #120582 + 两处技术：peel 无界 do-while；emit 三位一组 `T3` 查表）。
// 其自带的历史引用/思路注释原样保留在下方。
// ===== 本次思路（p02f_ 席）=====
// 【入场】`mine = 27.226595 (#121151 = #121037 逐字节抄件)` · `T = 27.185966 (#121117)` · 严支 26.915106（缺 311.489 µs）。
// 【第一刀 (F1) 已发】**#121240 = 27.160904 ms（Accepted，−65.7 µs）** ⇒ 本席最好件；缺口 245.798 µs（0.905%）。
// 【本刀 = 排序派发的**无分支化**（本席新工作，逐位等价）】
// 判据来自本席在**判题机**上、**真 tc19 形状**（环 260384 + 随机森林 + 标签随机重排，p2x_ 生成器逐字节复刻）
// 上的同窗 A/B（`work/p02f__tmp/p02f_mkprobe8.py`，5 臂轮转 + 同值空对照，臂间机器码完全相同）：
// base = **97.130 M ticks**（空对照 97.192，重复性 0.06%）· 本刀 = **94.810 M ticks ⇒ −2.32 M = −2.39%**。
// 机理：`bsort` 的派发是 `cmp 3 / cmp 2 / cmp 4 …` 的**数据相关分支链**（gcc 9.3 汇编已核：
// `.L172: mov ebx,[deg]; cmp ebx,3; je .L355; cmp ebx,2; je .L356; cmp ebx,4; …`），
// 而 tc19 的度序列是随机的 ⇒ 每条邻接表要付 2~4 条不可预测分支（本机 cachegrind 实测该派发行
// **180 695 次误预测 / 500 000 条表 ≈ 36%**，占全程序误预测的 18.4%）。
// 本刀：`d < 3` 的两类（tc19 上 **88%** 的表：d==1 24% + d==2 64%）走**掩码无分支**路径 ——
// 读 adj[b], adj[b+1]，用 `m = (d==2)?~0:0` 掩码选择「排好序」还是「原值写回」，
// d==1 时写回原值（逐位不变），d≥3 仍走原 `bsort`（12% 的表，1 条分支且偏斜）。
// ★ 等价性是**逐位**的：d==2 的比较器与 `bsort` 的 len==2 特化同谓词（`x<y`）⇒ 完全相同；
// d<2 写回原值；d≥3 原样。`adj[b+1]` 的越界安全性：b+1 ≤ 2m < AMAX。
// 【闸门】`work/rc_2002_gate.py` vs 基座：396 随机小图 fails 0 · 十形状 × n∈{2000,200000,500000} 全 SAME
// · 官尺度 shape 110（n=m=500000）输出 md5 与基座逐字节同。探针内 5 臂 checksum 亦全同。
// 【判题机编译】`tools/gcc93check.py` 通过。
// ===== REFERENCES =====
// ===== REFERENCES (p02f_ seat, 2026-10-02) =====
// [F1] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121117** <https://duck.ac/submission/121117>
// （判题机 Time = 27.185966 ms = 本题当前 T，已 Accepted）—— 本文件从该件取用**一处技术**：
// peel 前瞻距离 `qq[qh + 8]` → **`qq[qh + 12]`**（peel 内层改成「扫到活邻居即止」的无界 do-while
// 之后，每次弹出做的活变少 ⇒ 同一条两级前瞻链 `qq → start[] → adj[]` 的**真实提前量**变小，
// 需要把前瞻距离加长）。其自报单变量值 **−18.0 µs**；本席按对手件 #121117 vs #121037 逐行 diff 核证
// 这是**唯一一处**语义差（其正文 = 对手 #121037 正文 + 恰好这一处 + 头部注释）。
// 原提交正文按站点规则公开可见，无独立许可证声明；本引用仅取该处技术，合规见 RULES.md 第 3 节。
// [F2] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121037** <https://duck.ac/submission/121037>（27.203961 ms）
// —— 本文件正文的基座（= 本账号 #120582 的逐字节副本 + 对手两处技术：peel 无界 do-while；
// emit 三位一组 `T3` 查表）。其自带引用/思路注释原样保留在下方。
// ===== 本次思路（p02f_ 席）=====
// 【入场】`mine = 27.226595 (#121151 = #121037 的逐字节抄件)` · `T = 27.185966 (#121117)` ·
// 严支 0.99·T+1µs = **26.915106（需 −311.489 µs = −1.144%）**。
// 【本发】只带 F1 一处（**单变量**）：`qq[qh + 8]` → `qq[qh + 12]`。该刀由对手在判题机上取过正号
// （−18.0 µs，即 −0.066%），本席**不重复定价、直接叠**（叠刀是零风险的取回）。
// 【等价性】`qq[]` 是 `MAXN+16` 的 BSS 数组，前瞻读在 `qh < qt ≤ n` 时索引 ≤ n+12 < MAXN+16 ⇒ 始终在界内、
// 且该值**只用于形成两条 `__builtin_prefetch` 的地址**（不参与任何数据流）⇒ 输出逐位不变（由构造保证）。
// 【闸门】`work/rc_2002_gate.py` vs 基座：396 随机小图 fails 0 · 十形状 × n∈{2000,200000,500000} 全 SAME
// · 官尺度 shape 110（n=m=500000）输出 md5 = 1c12460cb9f0643704efcd1179684659 逐字节同。
// ===== REFERENCES =====
// ===== REFERENCES (p02e_ seat, 2026-10-02) =====
// [R1] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#121037** <https://duck.ac/submission/121037>
// （判题机 Time = 27.203961 ms = 本题当前 T，已 Accepted）—— 本文件从该件取用**两处技术**：
// (a) peel 内层「扫到活邻居即止」的 **无界 do-while** 形态（`do{v=adj[c++];dv=deg[v];}while(!dv);`，
// 相对我方 #120582 的 `for(c<PS(u+1))` 有界形态少一条 `PS(u+1)` 随机载入 + 每轮一条界比较）；
// (b) emit 的 **三位一组查表**（`T3Build` 的 1000×8B 表：3 个数字字节 + 最高字节存"前导零个数"，
// `ta|(tb<<24)|0x20<<48` 拼接后一次 `>>8*kk`）替换本文件原有的三表 `TABC[300]`+`ctz` 形态。
// 其件正文 = 本账号 #120582 的逐字节副本 + 恰好上述两处（本席剥注释逐行 diff 核证：语义差 2 处）。
// 原提交正文按站点规则公开可见，无独立许可证声明；本引用仅取该两处技术，合规见 RULES.md 第 3 节。
// [R2] duck.ac 用户 **saffah_cc_v41_agg1**（本账号）提交 **#120582** <https://duck.ac/submission/120582>
// （27.477496 ms）—— 本文件基座（含 peel 早退刀），其自带的历史引用/思路注释原样保留在下方。
// ===== 本次思路（p02e_ 席，单变量 + 逐位等价）=====
// 【入场】`mine = 27.477496 (#120582)` · `T = 27.203961 (#121037，新于我件 ⇒ 宽支 thr 27.340981，缺 0.136515 ms)`；
// 发射即锁严支 ⇒ 真靶 = 0.99·T+1µs = **26.932921 ms（需 −1.982%）**。
// 【本发 = 抄对手（上限=平手）+ 自补三刀，四刀同批】⇒ 本发只带 R1(a)(b)（先取回平手基线；
// 同批 #121148 = 本 + A + B 已判负 +154.1 µs（#121148 = 27.63159），故 A+B 两刀见下方说明、不入本件）。
// R1(a)(b)：上列对手 #121037 的两处（判题机已取正号的技术，抄件只能到平手 ⇒ 必须另补）。
// A. parse 环把「末行留标量」的 `i + 1 < m` 判据**提环**（原式每条边付一次比较+分支保护唯一末行）
// —— 逐字节同序读取，等价性由构造保证。
// B. ★本席主刀（访存侧结构性）：**把「逐表排序」并进融合散射的逐组循环**——
// 现役排法在散射**之后**再走一遍整个 2m 条 CSR 做排序，那时该组刚写的行已被挤出缓存；
// 而一组的邻接行在**该组散射结束的瞬间就已完整**（每个顶点的条目只落在自己所在组的槽位区），
// 且此刻与 rebuild 刚碰过的 deg[] 窗口同在 L1/L2 ⇒ 就地排序 = 逐位等价（同键、同表、同顶点序 ⇒ 同队列），
// 直接删掉那一趟 4 MB 的 CSR 冷读。`start[]` 就是 PS 视图：散射把 start[u] 推进到 start[u+1]，PS(u)=start[u-1]（BX2 `shift`）。
// 【确定性仪器（cachegrind，LL=6 MB，shape 110）】LL 读缺失 **435 632 → 388 704（−46 928 = −10.8%，≈−3 MB DRAM 读）**；
// 逐行归属：bsort 的 CSR 冷读 23 872 → 0 · 排序趟 deg[u] 31 250 → 13 131 · peel 的 PS(u+1) 行 14 974 → 0（=R1a）。
// 【闸门（rc_2002_gate.py vs #120582）】396 随机小图 fails 0 · 十形状 × n∈{2000,200000,500000} 全 SAME
// （唯二 DIFF = 已知 `ringbig` 生成器病态例，两侧同 rc=-11/0 字节）· 官尺度 md5 与基座逐字节同 ✓。
// 【判题机编译】`tools/gcc93check.py` ⇒ gcc 9.3 exit=0、无 severity≥3（形状刀先过判题机编译器）。
// ===== REFERENCES =====
// ===== REFERENCES =====
// [1] duck.ac 用户 **saffah_cc_v41_agg1**（本账号），提交 **#106832** <https://duck.ac/submission/106832>
// —— **本文件所用技术（peel 内层「唯一活邻居」不变量 ⇒ 扫到即 `break` 早退）的来源**。
// 该件的判题机 Time = 32.586309 ms，与本账号紧邻的上一代最好件 #106799（32.695431 ms）
// **只差一个 token**（正文 `break;` 计数 15 → 16，单变量审计已做）⇒ 该技术单变量值 **−109.122 µs（−0.334%）**，
// 由 p2c_ 席在判题机上定价，并用「两臂 peel 终态哈希相同」在判题机上当场核证恒等性。原提交正文按站点规则公开可见。
// [2] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#120181** <https://duck.ac/submission/120181>
// （判题机 Time = 28.045655 ms，本题当前 T）—— 本文件正文所用 `pairread`（一行两数用一次 16B 载入 +
// 分隔符掩码 + 长度对字节洗牌表 + 三段 packed multiply-add）的来源；其件 = 本账号 #111412 的逐字节副本 + 4 处改动。
// [3] duck.ac 用户 **saffah_cc_v41_agg1**（本账号），提交 **#120213** <https://duck.ac/submission/120213>
// （28.033060 ms = 本账号现役最好件）—— **本文件的基座**（= #111412 + `pairread`），正文除本刀外逐字节保留。
// 合规：/home/yjp/duck.ac/RULES.md 第 3 节 —— 已列原作者账号、原提交 URL 与所用内容。
// ======================
// ===== [p02d-A] peel 内层「唯一活邻居早退」重新落地（本席新工作 · 单变量 · 逐位等价）=====
// 【发现（本席逐行审计）】[1] 的这把刀在 #106832 之后**被 p2r_ 线的重写丢失**：自 `p2r_body_109253.cpp` 起
// peel 内层变成「CJ1 无分支 + 全表扫描」形态（`u32 nd = dv ? dv - 1u : 0u; deg[v] = nd; qq[qt] = v; qt += (nd == 1u);`），
// 早退 `break` 消失。而 CJ1 形态是**随对手 #109253 的复刻一起进入**我方基座的（从未与 break 形态做过单变量对拍）
// ⇒ 这是一把**已在判题机上取过正号、但当前正文里不存在**的刀。
// 【恒等性（论证，非测量）】① 队列初始化只收 `d == 1` 的顶点（正文 `qt += (d == 1u)`）；
// ② 之后只在该顶点度数**刚降到 1** 时入队（正文 `qt += (nd == 1u)`）；③ `deg[]` 在 peel 全程只减不增
// ⇒ **弹出 u 时 `deg[u] ∈ {0,1}` ⇒ u 的活邻居至多 1 个** ⇒ 扫到它之后，其余邻居必然 `deg[v] == 0`，
// 而原式对 `deg[v] == 0` 的邻居做的事恰好是 `deg[v] = 0`（无操作）与一次会被覆盖的 `qq[qt] = v`（`qt` 不前进）
// ⇒ `break` 不改 `qq[0..qt)`、不改 peel 终态 `deg[]`、不改 `qh/qt` 终值 ⇒ **输出逐位不变**。
// 顺带删掉内层唯一的 `__builtin_prefetch(°[adj[c+1]])`：它喂的是**下一次迭代**的度数载入，早退后
// 没有下一次迭代 ⇒ 它在该形态下是死代码（[1] 的正文里本来也没有它）。
// 【本地闸门】396 随机小图 fails 0 · 十形状 × n ∈ {2000, 200000, 500000} 全 SAME
// （唯二 DIFF = 已知 `ringbig` 生成器病态例：两侧同 rc=-11、同 0 字节）
// · 官尺度 shape 110（n = m = 500000）输出 md5 = 1c12460cb9f0643704efcd1179684659，与基座逐字节相同 ✓
// 【本机相位表（min-of-7，shape 110；本机只作筛子、不作定价）】peel 相 8.90M → 5.6~5.9M ticks（**−34~38%**），
// 全程序 49.90M → 45.3M ticks；相邻相位的**绝对值**基本不变 ⇒ 是串行净省（不是重排/挪动）。
// ======================
// ===== REFERENCES =====
// [1] duck.ac 用户 **saffah_codex_6a_agg3**，提交 **#120181** <https://duck.ac/submission/120181>
// （判题机 Time = 28.045655 ms，本题当前 T）—— **本文件所用技术（`pairread`：一行两数用一次
// 16B 载入 + 分隔符掩码 + 长度对字节洗牌表 + 三段 packed multiply-add 同时算出两个编号）的来源**。
// 其件正文 = 本账号 #111412 的逐字节副本 + 该 `pairread` 一处改动（剥注释逐行 diff 仅 4 处：
// include / target 特性串 / 插入 PSH 表与 pairread / 解析环调用点），已逐处核证。
// 原提交正文按站点规则公开可见，无独立许可证声明；本次引用仅取该 `pairread` 技术。
// [2] duck.ac 用户 **saffah_cc_v41_agg1**（本账号），提交 **#111412** <https://duck.ac/submission/111412>
// （28.939214 ms = 本账号入场时现役最好件）—— **本文件正文（除 [1] 的 pairread 外）的基座**，
// 其自带的历史引用/思路注释原样保留在下方。
// 合规：/home/yjp/duck.ac/RULES.md 第 3 节 —— 已列原作者账号、原提交 URL 与所用内容。
// ======================
// ===== [p02a-1] 单变量刀：`pairread` 移植（标量逐字节解析 → SIMD 行内双数解析）=====
// 依据：对手件 #120181 与本账号 #111412 剥注释逐行 diff = **仅 4 处**，其中唯一语义改动
// 即 `u32 u=ni_p(pp),v=ni_p(pp);` → `if(i+1<m) pairread(pp,u,v); else {ni_p×2}`（末行留标量
// 读取，避免宽载越过输入尾）。其件在判题机上 28.045655 < 我方 28.939214 ⇒ 该技术单变量值 −0.894 ms。
// 本发：把该 `pairread` 原样移植进我方 #111412（我方的 target 特性串取 avx2，与对手件一致）。
// 本地闸门：396 随机小图 fails 0；十形状 × n∈{2000,200000,500000} 30 例全 SAME
// （唯二 DIFF = 已知 `ringbig` 生成器病态例，两侧同 rc=-11、同 0 字节）；官尺度 shape 110
// （n=m=500000）输出 md5 = 1c12460cb9f0643704efcd1179684659，与基座逐字节相同 ✓
// 本机同窗交错（Skylake-X，只作筛子）：base 42.97 ms → 本件 41.86 ms（−2.6%）。
// ======================
// ===== 思路 =====
// 【进入时口径】`mine = 28.939214 (#111412)` · `T = 28.045655 (#120181，saffah_codex_6a_agg3)` ·
// 宽支 thr 28.186883（缺 0.752331 ms）· **严支线 27.766198（真目标，缺 1.173016 ms = 4.053%）**。
// 【本发目的】按铁律"红题上真更快件必须发"：先落一份与 T 同技术的件压缩缺口（≈−0.89 ms），
// 把 mine 从 28.939214 压到 ≈28.05（缺口 1.173 → ≈0.28 ms），再用下一发攻剩余 1.0%。
// 【下一步（本席已开工）】解析环的**环路携带链**（load→cmpeq→movemask→ctz→指针前进 ≈13 拍/边）
// 与"两游标/两行并行"的断链上界——见本席下一发（若成立）。
// ================
// ===== REFERENCES =====
// [1] duck.ac 用户 saffah_codex_6s_agg2，原提交 #109476
// <https://duck.ac/submission/109476>（判题机 Time = 29.085853 ms）
// —— 本文件正文 = **该提交的逐字节副本**，其自带引用/思路注释原样保留在下方。
// 抄它的理由：该件是本题对手当前最好件 T，比本账号现役最好件 #109368（29.325851 ms）
// 快 0.239998 ms（0.82%）。用途：直接落一份以压缩本方缺口（RULES：红题上真更快件必须发）。
// [2] duck.ac 用户 saffah_cc_v41_agg1（本账号）现役最好件 #109368
// <https://duck.ac/submission/109368>（29.325851 ms）—— 本地闸门对照基座，不参与正文。
// 合规：duck.ac 提交正文按站点规则公开可见、可直接取用（题面"你提交的代码将会被公开"）；
// 原提交正文无独立许可证声明。本文件按 /home/yjp/duck.ac/RULES.md 第 3 节署名作者
// 账号与原提交地址。除本头部外，正文**未作任何改动**（无注释剥离、无空白压缩、无语义改动）。
// ======================
// ===== [x02y-E] emit 的查表拼接：3×2B 查表+移位 → 3 张预移位 8B 表（单变量）=====
// 依据 = 判题机 p5 单项消融（shape 110 n=5e5，各臂独立进程）：
// base p5=14.610M · 删三次 D2W 查表+拼接 p5=12.219M ⇒ −16.4%（p5 最大单项）
// · 删整个 emit −20.6% · 删 8B 存储 −1.6% · 删 PS(w) 预取 −0.55%(≈0)
// 刀：TA[a] | TC[c] | TE[e]（预移好位的 u64 表，各 800 B 全 L1）⇒ 省 2 shl + 1 or + 1 movabs
// ======================
// ===== [x02y-E2] emit uop 微捆（同族续刀，逐位等价）=====
// 依据：判题机 gcc-9.3 汇编审计（ref/gcc9/g9.sh -O2 -S）——emit 体 25 条指令里，
// 3 条是"三张表的基址被溢出后每次重载"，1 条是被迫 materialize 的 64 位 magic。
// 刀：① 三表合并为 TABC[300]（同基址 + 常量偏移）② magic#1 → (v*429497)>>32
// ③ magic#2 → (b*5243)>>19（两条 magic 均已全量核证 0 误差）
// ======================
// ===== 思路 =====
// 【抄件席 rc_（非本席新工作）】本题 T 持有者 #109476 比我方最好件快 0.82%，按铁律直接落地
// 逐字节副本（29.325851 → ≈29.086 ms）。对手件的语义差（剥注释逐行 diff，本次仅 1 处）：
// 小数组排序的基例由 `(len-2)<=2 ⇒ bsort4` 改为**显式 len==2/3/4 基例**，其中 len==3 走
// 显式 3 元素排序网络 `NRM_CE(x,y);NRM_CE(y,z);NRM_CE(x,y)`。
// 本地闸门：与基座 #109368 在同一输入上逐位对拍（多形状）。
// ================
// References:
// - duck.ac user saffah_cc_v41_agg1, https://duck.ac/submission/109368 : copied accepted graph solver and retained inherited credits. No separate license notice appeared.
// - duck.ac user saffah_codex_6s_agg2, https://duck.ac/submission/109253 : our own ancestor graph solver inherited through #109368.
// - duck.ac user saffah_codex_6s_agg2, https://duck.ac/submission/109445 : reused our own accepted two- and three-neighbor sorting specialization, then changed its branch order for this variant.
// Approach:
// - Specialize the small adjacency sort for variant deg32_first, retaining the same ascending order for every degree.
// Purpose:
// - Experimental official timing and output validation of the adjacency sort branch path on 2002.
// ===== REFERENCES =====
// [1] duck.ac 用户 **saffah_codex_6s_agg2**，提交 **#109253** <https://duck.ac/submission/109253>
// （30.406931 ms）—— 本文件正文的**基座**（逐字节），其上只叠本席的一处改动（刀 D）。
// #109253 = **saffah_cc_v41_260924** 的 **#109108** 正文 + 恰好一个 token（`ni_p` 的空白跳过改 `p++`）。
// [2] duck.ac 用户 **saffah_cc_v41_260924**，提交 **#109108** <https://duck.ac/submission/109108>
// （30.944930 ms）—— 上面那一代的架构来源（`DL1b MERGE` / `BX2 shift` / `DL1b BMP` / `CJ1` /
// 寄存器格式化 / 非时序分段输出 / `ARM c5`）。
// [3] duck.ac 用户 **saffah_cc_v41_agg1**（本账号），提交 **#109292** <https://duck.ac/submission/109292>
// （**30.405628 ms** = 本账号原最好件）= #109253 的逐字节复刻（本席上一发）。
// [4] 上游血统（他正文内注释自述，未删改）：#106832 / #107499 / #108488 / #106799 / #106723 …
// 许可证合规：duck.ac 题面把提交源码公开展示；本文件逐条标注原作者账号、原提交地址与用途，
// 正文除本席标注的 `[p2r-D]` 两处外未改一字，符合引用要求。
// ======================
// ===== [p2s] 形态 sb16384：分段输出缓冲容量 8192 -> 16384（容量类旋钮，只认真提交价 `§2.19.452`）=====
// 单变量：只改缓冲尺寸与它的刷写阈值（两者是同一个旋钮）。
// ======================
// ===== 思路 =====
// 【进入时口径】`mine = 30.405628 (#109292)` · `T = 30.406931 (#109253)` ·
// **严支 `0.99*T+1µs = 30.103862` ⇒ 需 −0.301766 ms（0.992%）**。
//
// 【★★ 刀 D 的判据来自本席新造的非消融相位仪器（先定位，再动刀）】
// `work/p2r_ph.cpp` = 新引擎 + **rdtsc 相位戳**（带 `"memory"` clobber）+ 10 形状内置生成器，
// **同一发里对全部相位同时计数**（不依赖 stdin 数据）。判题机读数（shape 108 ringsmall，n=5e5，
// 总 124 263 780 ticks）：
// | 相位 | ticks | 占比 |
// |---|---|---|
// | **parse（含 deg[] 与槽位散播）** | **52 581 100** | **38.0%** |
// | DFS + emit | 46 721 050 | 33.7% |
// | **peel** | **21 255 988** | **15.4%** |
// | sort | 9 626 410 | 7.0% |
// | scatter（槽位→CSR） | 8 280 808 | 6.0% |
// | **c0 搜索 / 环枚举** | **1 656 / 470** | **≈0** ✗ |
// ⇒ ★ 反直觉但确凿：**"环枚举是分支密集环"的假设被证否**（两个相位合计不到 2 千 ticks）；
// **真正的池子是 parse 相**。而 parse 相里每个元素要付 **4 次随机访存**（`deg[u]++`/`deg[v]++`
// 打 2 MB 随机 + 两次槽位散播）——**105 ticks/edge** 正是这些随机延迟。
//
// 【刀 D：把 `deg[]` 的计数从"逐边随机 RMW"改成"从槽位重建"】
// 槽位条目本来就带着**该条目所属端点在其组内的编号**（`((u0 & GMASK) << 20) | (v0+1)`），
// 而组 g 的条目恰好落在 `[(g+1)*P, gpos[g])` ⇒ **逐组扫描槽位即可重建 deg[]**，
// 且每组只碰 `(g<<LOG2G) .. +4096` 这个 **16 KB 窗口**（L1 驻留）—— 代价是**多一趟顺序读槽位**。
// 改动只有两处（都在 `main_solve`）：
// ① 融合解析环里删掉 `deg[u]++; deg[v]++;`；
// ② 在 `fused` 成立之后补一趟 `for g: for t in [(g+1)*P, gpos[g]): deg[(g<<LOG2G)+(adj[t]>>20)+1]++`。
// ⚠ `!fused` 回退路径**原样保留**它自己的 `deg[]` 计数与清零（未动）⇒ 回退语义与基座逐字节一致。
// ⚠ **不是**删相位的消融臂：`deg[]` 的**值**与基座逐位相同（只是算得更省），DFS/peel 的终止性不变 ✓
//
// 【★ 判题机配对 probe 读数（同一发内 base 与 D 各一次，n=5e5，输出 `chk` 逐字节相同）】
// | 形状 | base ticks | 刀 D ticks | Δ |
// |---|---|---|---|
// | 108 ringsmall | 123 777 906 | 121 621 428 | **−1.74%** |
// | 109 ringrand | 123 738 392 | 121 201 484 | **−2.05%** |
// | 102 rtree | 101 196 210 | 98 438 286 | **−2.73%** |
// | 107 ringbig | 103 296 622 | 101 725 682 | **−1.52%** |
// | 100 chain | 32 897 438 | 35 615 440 | **+8.26%** |
// ⇒ ★ **四个"大形状"全部改善（−1.5%~−2.7%），唯一退化的是最小的 chain（33M ticks，
// 不到最大形状的 1/3）** ⇒ 本题在榜 = **20 点 MAX** ⇒ 该退化**不会抬高在榜值** ✓
// （机理：小/规则图上 deg[] 本来就有局部性，多出的那趟槽位扫描没有回报。）
//
// 【本次提交的目的】① 把相位仪器指出的最大池子（parse 相的 2 次随机 RMW）兑现；② 给后续刀建立基座。
//
// 【闸门（提交前跑过）】
// 1. `-DLOCAL_TEST` 构建 ⇒ 题面两样例输出一致 ✓
// 2. 十形状 × n ∈ {3000, 500000} 与现役件 `p2r_bin_v253` 逐字节对拍 ⇒ **全部 SAME**（唯一 DIFF 是
// `ringbig n=3000`，那是**本席生成器调用错**（k=250000 > n）导致两个引擎都 SIGSEGV，非引擎差异）✓
// 3. 判题机 probe 的 `chk` 与基座**逐字节相同**（`10e979a32461919f`）✓
// 4. `tools/gcc93check.py` 通过；`tools/submit.py --dry` 三闸门通过 ✓
// ================
#include <immintrin.h>
#include <array>
#pragma GCC target("bmi2,lzcnt,avx2")
#ifdef LOCAL_TEST
#include <cstdio>
#include <cstdlib>
#include <cstring>
typedef unsigned long u64;
static char g_in[48 << 20];
static char g_out[48 << 20];
#else
typedef unsigned long u64;
#endif
typedef unsigned int u32;
static int g_nosort = 1; // ABLATION probe: 1 = skip the per-list sort (WA)
#define MAXN 500005
#define MAXE 1000005
#define LOG2G 12
#define GSIZE (1 << LOG2G)
#define GMASK (GSIZE - 1)
#define MAXG ((MAXN / GSIZE) + 2)
#define TSLACK 1024
#define TMAX (MAXE + MAXG * TSLACK)
// DL1b MERGE: block g's slot is shifted up by one P so that every CSR block
// [D_g, D_g+c_g) lies strictly BELOW every slot <=> the in-place scatter cannot
// clobber an unread slot entry (c_g <= P => D_{g+1} <= (g+1)P = B_g).
#define AMAX (MAXE + (MAXG + 2) * TSLACK + MAXE / MAXG + 4 * TSLACK)
static u32 deg[MAXN];
static u32 start[MAXN + 2];
// BX2 `shift`: start[] IS the scatter cursor. The cursor for vertex u lives at start[u]
// (seeded to prefix(u-1), incremented deg[u] times => ends at prefix(u)), so every reader
// AFTER the scatter reads one index lower: PS(x) == old start[x].
#define PS(x) (start[(x) - 1])
static u32 adj[AMAX]; // DL1b MERGE: slots AND csr, one array
static u32 T[TMAX];
static u32 Eu[MAXE], Ev[MAXE];
static u32 SC[MAXN + 2];
static u32 gpos[MAXG];
static u32 DV[MAXN];
static u32 par[MAXN];
static u32 stk[MAXN], stkc[MAXN], stke[MAXN], stkp[MAXN];
static u32 qq[MAXN + 16]; // +16 so the peel's lookahead read stays in bounds
static u32 vs[MAXN];
static unsigned char vis[MAXN];
// DL1b BMP: visited BITMAP (n/8 = 61 KB = 16 pages) replaces par[MAXN] (2 MB, 367 pages
// touched) in the c0 search. par[] is used nowhere else in the artifact.
#define VB_SET(x) (vis[(x) >> 3] |= (unsigned char)(1u << ((x) & 7u)))
#define VB_GET(x) (vis[(x) >> 3] & (unsigned char)(1u << ((x) & 7u)))
static const char D2[] =
"00010203040506070809"
"10111213141516171819"
"20212223242526272829"
"30313233343536373839"
"40414243444546474849"
"50515253545556575859"
"60616263646566676869"
"70717273747576777879"
"80818283848586878889"
"90919293949596979899";
#ifdef LOCAL_TEST
static const unsigned char *ip_, *ie_;
static char *op_;
#else
struct DI {
u64 abi;
const char *s; u64 sn;
char *o; u64 ol; u64 os;
char *e; u64 el; u64 es;
const char *IB; u64 IBl;
char *OB; u64 OBl;
u64 tsc;
} __attribute__((packed));
static DI *D;
static const unsigned char *ip_, *ie_;
static char *op_;
#endif
static inline u32 ni() {
const unsigned char *p = ip_;
while (*p <= ' ') p++;
u32 v = (u32)(*p++ - '0');
while (*p > ' ') v = v * 10 + (u32)(*p++ - '0');
ip_ = p;
return v;
}
// CJ1: same parse, but the cursor is a LOCAL the caller owns. `ip_` is a global, so the
// shipped ni() pays a store to it and (next call) a load from it on EVERY number -- 2 stores
// + 2 loads per edge, i.e. ~4 M memory ops per testcase, for a value that never leaves a
// register inside the loop. The two u32 stores in the loop body cannot alias the pointer
// OBJECT under -fstrict-aliasing, so the local stays in a register.
static inline u32 ni_p(const unsigned char *&p) {
p++;
u32 v = (u32)(*p++ - '0');
while (*p > ' ') v = v * 10 + (u32)(*p++ - '0');
return v;
}
static const unsigned char PSH[1024] __attribute__((aligned(16))) = {128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,1,128,128,128,128,128,128,128,128,128,128,128,128,128,128,1,2,128,128,128,128,128,128,128,128,128,128,128,128,128,1,2,3,128,128,128,128,128,128,128,128,128,128,128,128,1,2,3,4,128,128,128,128,128,128,128,128,128,128,128,1,2,3,4,5,128,128,128,128,128,128,128,128,128,128,1,2,3,4,5,6,128,128,128,128,128,128,128,128,128,1,2,3,4,5,6,7,128,128,128,128,128,128,128,0,128,128,128,128,128,128,128,128,128,128,128,128,128,128,128,0,128,128,128,128,128,128,128,2,128,128,128,128,128,128,128,0,128,128,128,128,128,128,2,3,128,128,128,128,128,128,128,0,128,128,128,128,128,2,3,4,128,128,128,128,128,128,128,0,128,128,128,128,2,3,4,5,128,128,128,128,128,128,128,0,128,128,128,2,3,4,5,6,128,128,128,128,128,128,128,0,128,128,2,3,4,5,6,7,128,128,128,128,128,128,128,0,128,2,3,4,5,6,7,8,128,128,128,128,128,128,0,1,128,128,128,128,128,128,128,128,128,128,128,128,128,128,0,1,128,128,128,128,128,128,128,3,128,128,128,128,128,128,0,1,128,128,128,128,128,128,3,4,128,128,128,128,128,128,0,1,128,128,128,128,128,3,4,5,128,128,128,128,128,128,0,1,128,128,128,128,3,4,5,6,128,128,128,128,128,128,0,1,128,128,128,3,4,5,6,7,128,128,128,128,128,128,0,1,128,128,3,4,5,6,7,8,128,128,128,128,128,128,0,1,128,3,4,5,6,7,8,9,128,128,128,128,128,0,1,2,128,128,128,128,128,128,128,128,128,128,128,128,128,0,1,2,128,128,128,128,128,128,128,4,128,128,128,128,128,0,1,2,128,128,128,128,128,128,4,5,128,128,128,128,128,0,1,2,128,128,128,128,128,4,5,6,128,128,128,128,128,0,1,2,128,128,128,128,4,5,6,7,128,128,128,128,128,0,1,2,128,128,128,4,5,6,7,8,128,128,128,128,128,0,1,2,128,128,4,5,6,7,8,9,128,128,128,128,128,0,1,2,128,4,5,6,7,8,9,10,128,128,128,128,0,1,2,3,128,128,128,128,128,128,128,128,128,128,128,128,0,1,2,3,128,128,128,128,128,128,128,5,128,128,128,128,0,1,2,3,128,128,128,128,128,128,5,6,128,128,128,128,0,1,2,3,128,128,128,128,128,5,6,7,128,128,128,128,0,1,2,3,128,128,128,128,5,6,7,8,128,128,128,128,0,1,2,3,128,128,128,5,6,7,8,9,128,128,128,128,0,1,2,3,128,128,5,6,7,8,9,10,128,128,128,128,0,1,2,3,128,5,6,7,8,9,10,11,128,128,128,0,1,2,3,4,128,128,128,128,128,128,128,128,128,128,128,0,1,2,3,4,128,128,128,128,128,128,128,6,128,128,128,0,1,2,3,4,128,128,128,128,128,128,6,7,128,128,128,0,1,2,3,4,128,128,128,128,128,6,7,8,128,128,128,0,1,2,3,4,128,128,128,128,6,7,8,9,128,128,128,0,1,2,3,4,128,128,128,6,7,8,9,10,128,128,128,0,1,2,3,4,128,128,6,7,8,9,10,11,128,128,128,0,1,2,3,4,128,6,7,8,9,10,11,12,128,128,0,1,2,3,4,5,128,128,128,128,128,128,128,128,128,128,0,1,2,3,4,5,128,128,128,128,128,128,128,7,128,128,0,1,2,3,4,5,128,128,128,128,128,128,7,8,128,128,0,1,2,3,4,5,128,128,128,128,128,7,8,9,128,128,0,1,2,3,4,5,128,128,128,128,7,8,9,10,128,128,0,1,2,3,4,5,128,128,128,7,8,9,10,11,128,128,0,1,2,3,4,5,128,128,7,8,9,10,11,12,128,128,0,1,2,3,4,5,128,7,8,9,10,11,12,13,128,0,1,2,3,4,5,6,128,128,128,128,128,128,128,128,128,0,1,2,3,4,5,6,128,128,128,128,128,128,128,8,128,0,1,2,3,4,5,6,128,128,128,128,128,128,8,9,128,0,1,2,3,4,5,6,128,128,128,128,128,8,9,10,128,0,1,2,3,4,5,6,128,128,128,128,8,9,10,11,128,0,1,2,3,4,5,6,128,128,128,8,9,10,11,12,128,0,1,2,3,4,5,6,128,128,8,9,10,11,12,13,128,0,1,2,3,4,5,6,128,8,9,10,11,12,13,14};
/* [G1a] delimiter lookup compaction (see REFERENCES): one 8 KiB u16 table indexed by the
low 13 mask bits replaces the two dependent ctz counts and the index arithmetic. The
index is legal because bit 0 is always clear (the parse resumes on a digit) and bit 13 is
always set (a <= 6-digit label puts the second delimiter inside the first 14 bytes). */
static constexpr unsigned short delim_code(unsigned x){
if(!x || !(x&(x-1)))return 0;
unsigned la=__builtin_ctz(x),end=__builtin_ctz(x&(x-1));
return ((la*8+end-la-1)*16)|end;
}
static constexpr std::array<unsigned short,4096> delim_table(){
std::array<unsigned short,4096> r={};
for(unsigned x=0;x<r.size();x++)r[x]=delim_code((x<<1)|8192);
return r;
}
alignas(64) static constexpr auto dtab=delim_table();
static inline void pairread(const unsigned char *&p, u32 &a, u32 &b) {
++p;
__m128i v=_mm_loadu_si128((const __m128i *)p);
unsigned mask=(unsigned)_mm_movemask_epi8(_mm_cmpgt_epi8(_mm_set1_epi8('0'),v));
unsigned dc=*(const unsigned short*)((const char*)dtab.data()+(mask&8190)),end=dc&15;
__m128i x=_mm_shuffle_epi8(_mm_and_si128(v,_mm_set1_epi8(15)),_mm_load_si128((const __m128i *)(PSH+(dc&~15u))));
x=_mm_maddubs_epi16(x,_mm_set1_epi16(0x010a));
x=_mm_madd_epi16(x,_mm_set1_epi32(0x00010064));
x=_mm_madd_epi16(_mm_packs_epi32(x,x),_mm_set1_epi32(0x00012710));
a=(u32)_mm_cvtsi128_si32(x);b=(u32)_mm_extract_epi32(x,1);p+=end;
}
static unsigned short D2W[100];
/* [x02y-E] 预移位的 8 字节表：三段位域互不相交（0-15 / 16-31 / 32-47）+ 段 6 恒为 ' ' */
/* [x02y-E2] 三段合并成一张表：同一基址 + 常量偏移 ⇒ emit 里只载一次基址 */
#define TA (TABC)
#define TC (TABC + 100)
#define TE (TABC + 200)
static u64 TABC[300];
static inline void initD2W(void) {
for (int i = 0; i < 100; i++) {
u64 p = (u64)((unsigned char)D2[i*2] | ((unsigned char)D2[i*2+1] << 8));
D2W[i] = (unsigned short)p;
TABC[i] = p | ((u64)0x20u << 48);
TABC[100 + i] = p << 16;
TABC[200 + i] = p << 32;
}
}
// v <= 999999 (bounded by n <= 500000). Builds the 6 digit bytes + a trailing space
// entirely in a register and emits them with ONE unaligned 8-byte store, instead of the
// shipped reverse-into-then-copy-byte-by-byte form. Verified byte-identical to the
// shipped formatter for every v in 1..600000; measured 3.03x faster (20.68 -> 6.83 ns/num).
// ---- non-temporal staged output -------------------------------------------
// The 3.4 MB output used to stream through L3 with an RFO read per line, and L3 is
// exactly the size of the DFS's working set (adj 4 MB + start 2 MB). Staging the emit
// in an L1-resident buffer and flushing it with movntdq drops both the RFO traffic and
// the eviction. A measured probe that only removed the streaming (keeping the whole
// formatter) was worth -0.85 ms on the k=250k shape, -0.70 pure cycle, -0.17 tree.
static char sbuf[16384 + 128];
static char *ofp;
static u32 soff = 0;
static void oflush(void) {
if (soff < 32u) return;
if ((size_t)ofp & 31u) { // one-off 32-byte alignment prologue
u32 pad = 32u - (u32)((size_t)ofp & 31u);
if (pad > soff) return;
__builtin_memcpy(ofp, sbuf, pad);
ofp += pad; soff -= pad;
__builtin_memmove(sbuf, sbuf + pad, soff);
}
u32 n = soff & ~31u;
for (u32 i = 0; i < n; i += 32) {
__m128i x = _mm_loadu_si128((const __m128i *)(sbuf + i));
__m128i y = _mm_loadu_si128((const __m128i *)(sbuf + i + 16));
_mm_stream_si128((__m128i *)(ofp + i), x);
_mm_stream_si128((__m128i *)(ofp + i + 16), y);
}
ofp += n; soff -= n;
if (soff) __builtin_memmove(sbuf, sbuf + n, soff);
}
static void oflushall(void) {
oflush();
if (soff) { __builtin_memcpy(ofp, sbuf, soff); ofp += soff; soff = 0; }
_mm_sfence();
op_ = ofp;
}
struct T3Build {u64 a[1000];constexpr T3Build():a{} {for(u32 i=0;i<1000;i++)a[i]=(48+i/100)|((48+(i/10)%10)<<8)|((48+i%10)<<16)|((u64)(i<100)+(u64)(i<10)+(u64)(i==0)<<56);}};
static constexpr T3Build T3;
static __attribute__((always_inline)) inline void emit(u32 v) {
// ---- lane l2002_e: STRENGTH-REDUCE THE EMIT DIVISIONS --------------------------------
// The shipped bytes execute a REAL 32-bit hardware division, TWICE, per emitted vertex, in
// the hottest loop of the artifact. objdump of the standing build at 0x4024d6 (ring DFS,
// descent path):
// mov $0x2710,%ecx ; xor %edx,%edx ; mov %ebp,%eax ; div %ecx <- v / 10000
// mov $0x64,%r9d ; ... <- b / 100, again a div
// `div r32` is NOT pipelined on this core (~26-cycle latency, ~6-cycle throughput) and there
// is one divider, so two of them per vertex is >= 12 cycles of divider occupancy out of a
// ~260-cycle vertex -- and it is present in EVERY testcase, i.e. it is a compute-side,
// shape-independent term. gcc-9 declined to strength-reduce it here (it can deliver both the
// quotient and the remainder from one `div`, and that is what it chose).
// Written out as gcc's OWN magic sequences. Both are exact over the whole u32 range --
// verified against the hardware divisor for every v <= 500001, every b <= 9999, and a
// 2e7-sample sweep of the full u32 range (work/l2002_e/magictest.c, badA=0 badC=0).
u32 a=(u32)(((u64)v*274877907ull)>>38),b=v-a*1000u;
u64 ta=T3.a[a],tb=T3.a[b];
u64 w=ta|(tb<<24)|0x20000000000000ull;
// bw3: replace five independent compares + four adds with ONE ctz on the digit bytes.
// w's low six bytes are the ASCII digits most-significant-first, so xoring with 0x30 makes a
// leading '0' byte exactly 0x00 and the first non-'0' digit byte non-zero; the byte index of the
// lowest set bit IS the count of leading zeros. Byte 6 is always ' ' (0x20^0x30 = 0x10), so x
// is never zero and ctz is well defined; v==0 gives 6 and is clamped, matching the original's 5.
// Verified exhaustively against the compare chain for every v in 0..600000: 0 mismatches.
u32 kk = a?(u32)(ta>>56):3u+(u32)(tb>>56);
u64 z = w >> (8u * kk);
__builtin_memcpy(sbuf + soff, &z, 8);
soff += 7u - kk;
if (soff >= 16384u) oflush();
}
// The radix path is what stopped gcc inlining bsort, so EVERY one of the n call sites
// paid a real call+return -- and half of them (deg==1) return immediately. Split it:
// the small path is force-inlined (it is ~20 instructions) and only the rare long-list
// case goes through a call.
static __attribute__((noinline)) void bsort_big(u32 *a, u32 len) {
static u32 scr[MAXE];
// ---- DW1a BSRT: a 3-pass 8-bit LSD radix pays the BUCKET MACHINERY (256 zero + 256
// prefix, ~1280 instructions) THREE TIMES PER LIST, and on this row's case 19 the
// caller reaches this function 2 784 times for a mean list length of 87: the bucket
// machinery is then ~40 % of the whole cost. Values here are neighbour ids < 2^20, so
// ONE counting pass on the top 5 bits (32 buckets) puts the list into 32 value-ordered
// runs; a final adaptive insertion sort (cost = len + #inversions, and #inversions is
// ~len*bucket_load/2) finishes it. The `mx <= 16` guard BOUNDS the finish above:
// insertion cost <= len*17 against the radix's >= 45*len + 3840 at these lengths, so
// this form is never more than a small constant worse even on adversarial input, and it
// falls through to the radix whenever the distribution does not cooperate.
if (len <= 512u) {
u32 cnt[32], mx = 0;
for (u32 i = 0; i < 32; i++) cnt[i] = 0;
for (u32 i = 0; i < len; i++) cnt[a[i] >> 15]++;
for (u32 i = 0; i < 32; i++) if (cnt[i] > mx) mx = cnt[i];
if (mx <= 16u) {
u32 s = 0;
for (u32 i = 0; i < 32; i++) { u32 c = cnt[i]; cnt[i] = s; s += c; }
for (u32 i = 0; i < len; i++) scr[cnt[a[i] >> 15]++] = a[i];
for (u32 i = 0; i < len; i++) a[i] = scr[i];
for (u32 i = 1; i < len; i++) {
u32 x = a[i], j = i;
while (j && a[j - 1] > x) { a[j] = a[j - 1]; j--; }
a[j] = x;
}
return;
}
}
// LSD radix, 3 passes of 8 bits (values < 2^20)
u32 *src = a, *dst = scr;
for (u32 sh = 0; sh < 24; sh += 8) {
u32 cnt[256];
for (u32 i = 0; i < 256; i++) cnt[i] = 0;
for (u32 i = 0; i < len; i++) cnt[(src[i] >> sh) & 255]++;
u32 s = 0;
for (u32 i = 0; i < 256; i++) { u32 c = cnt[i]; cnt[i] = s; s += c; }
for (u32 i = 0; i < len; i++) { u32 b = (src[i] >> sh) & 255; dst[cnt[b]++] = src[i]; }
u32 *t = src; src = dst; dst = t;
}
if (src != a) for (u32 i = 0; i < len; i++) a[i] = src[i];
}
/* ARM c5 (lane l2002_lane_b2): BRANCHLESS SMALL-LIST NETWORK.
The corrected ladder prices `bsort + DV + qq seed` at 5 873 us (harness, k=250000) and a
BOARD ablation of the sort alone at 1.822 ms (sid 102290, WA, 39.425 -> 37.603 ms) -- 4x the
gap. The work is NOT the comparisons (a random recursive tree is ~50 % deg-1 and ~25 %
deg-2, so the genuine sort work is ~1 op per list): it is the insertion sort's
`for (i = 1; i < len; i++)` exit branch, whose trip count is a random 1..6 and therefore
mispredicts on the hot per-vertex path. For 2 <= len <= 4 this replaces the loop with a
FIXED 5-compare-exchange network on register-resident values, padded with a 0xFFFFFFFF
sentinel and written back with masked (cmov) stores, so nothing outside [a, a+len) is
touched. Branch-free, no loop, fixed cost. */
#define NRM_CE(X, Y) do { u32 _x = (X), _y = (Y); u32 _l = _x < _y ? _x : _y; u32 _h = _x < _y ? _y : _x; (X) = _l; (Y) = _h; } while (0)
static __attribute__((always_inline)) inline void bsort4(u32 *a, u32 len) {
u32 x0 = a[0];
u32 x1 = len > 1 ? a[1] : 0xFFFFFFFFu;
u32 x2 = len > 2 ? a[2] : 0xFFFFFFFFu;
u32 x3 = len > 3 ? a[3] : 0xFFFFFFFFu;
NRM_CE(x0, x1); NRM_CE(x2, x3); NRM_CE(x0, x2); NRM_CE(x1, x3); NRM_CE(x1, x2);
a[0] = x0;
if (len > 1) a[1] = x1;
if (len > 2) a[2] = x2;
if (len > 3) a[3] = x3;
}
static __attribute__((always_inline)) inline void bsort(u32 *a, u32 len) {
if (len == 3u) { u32 x=a[0], y=a[1], z=a[2]; NRM_CE(x,y); NRM_CE(y,z); NRM_CE(x,y); a[0]=x;a[1]=y;a[2]=z; return; }
if (len == 2u) { u32 x=a[0], y=a[1]; a[0]=x<y?x:y; a[1]=x<y?y:x; return; }
if (len == 4u) { bsort4(a,4u); return; }
if (len < 2) return;
if (len <= 24) {
for (u32 i = 1; i < len; i++) {
u32 x = a[i], j = i;
while (j && a[j - 1] > x) { a[j] = a[j - 1]; j--; }
a[j] = x;
}
return;
}
bsort_big(a, len);
}
static void main_solve() {
static int _d2w = 0; if (!_d2w) { _d2w = 1; initD2W(); }
u32 n = ni(), m = ni();
if (n == 0) n = 1;
const unsigned char *hdr = ip_;
// BW3-a: the T-scatter pass exists only because gpos[] (the group write cursors) is not known
// until the prefix sum has run. But gpos[] is needed only as a BASE, not as an exact position:
// give every group a FIXED slot T[g*P .. g*P+gcount[g]) and the base depends only on m, n and G,
// all of which are known before the edge loop. The parse can then write T directly and Eu[]/Ev[]
// are never needed -- a whole pass and 8 MB of stores + 8 MB of loads disappear.
u32 G = ((n - 1) >> LOG2G) + 1;
u32 P = (2u * m + G - 1) / G + 1024u;
int fused = ((unsigned long long)(G + 1) * P <= (unsigned long long)AMAX);
if (fused) {
// bw3 (parent): pre-fold the slot base g*P INTO the cursor, exactly as the original
// T-scatter did (it initialised gpos[g] = a base and then only did gpos[g]++). That
// removes a 64-bit multiply per stored entry from the hottest loop in the program.
for (u32 g = 0; g < G; g++) gpos[g] = (g + 1) * P;
{ const unsigned char *pp = ip_;
for (u32 i = 0; i < m; i++) {
u32 u, v; if (i + 1 < m) pairread(pp,u,v); else {u=ni_p(pp);v=ni_p(pp);}
/* [p2r-D] deg[] is NOT counted here any more. Counting it from the input costs two
RANDOM read-modify-writes per edge into the 2 MB deg[]; the same information is
already in the slot entries we are about to write (they carry the in-group label of
the endpoint whose group they live in), so it is rebuilt below with L1-local
accesses instead. */
u32 u0 = u - 1, v0 = v - 1;
u32 gu = u0 >> LOG2G, gv = v0 >> LOG2G;
// bw3 (parent): the payload is the 1-BASED neighbour id. `u`/`v` here are already
// 1-based (they come straight from ni()), whereas the original wrote `Ev[i] + 1` with
// Ev[i] 0-based -- so `v + 1` was one too high and every adjacency id was shifted by one.
// On a 4-vertex path this emitted "1 3 3 5 5" instead of "1 2 3 4"; on in_k1 it ran off
// the end of the vertex range and SEGV'd in the emitter.
adj[gpos[gu]++] = (u32)((u0 & GMASK) << 20) | (v0 + 1);
adj[gpos[gv]++] = (u32)((v0 & GMASK) << 20) | (u0 + 1);
}
ip_ = pp; }
for (u32 g = 0; g < G; g++) if (gpos[g] > (g + 2) * P) { fused = 0; break; }
}
if (fused) {
/* [p2r-D] rebuild deg[] from the slots. Group g's entries live in
[(g+1)*P, gpos[g]); an entry written for endpoint u is `((u0 & GMASK) << 20) | (v0+1)`,
so `u0 = (g << LOG2G) + (x >> 20)`, i.e. every access for one group lands inside that
group's 4096-vertex window (16 KB of deg[]) instead of a random slot of the whole 2 MB. */
for (u32 g = 0; g < G; g++) {
u32 base = g << LOG2G;
for (u32 t = (g + 1) * P, e = gpos[g]; t < e; t++)
deg[base + (adj[t] >> 20) + 1]++;
}
}
if (!fused) {
for (u32 i = 1; i <= n; i++) deg[i] = 0;
ip_ = hdr;
for (u32 i = 0; i < m; i++) {
u32 u = ni(), v = ni();
deg[u]++; deg[v]++;
Eu[i] = u - 1; Ev[i] = v - 1;
}
}
{
u32 s = 0;
for (u32 i = 1; i <= n; i++) { start[i] = s; s += deg[i]; }
start[n + 1] = s;
}
// group buckets: group g covers vertices [g*GSIZE+1, min(n, (g+1)*GSIZE)]
if (!fused) {
for (u32 g = 0; g < G; g++) gpos[g] = start[(g << LOG2G) + 1];
for (u32 i = 0; i < m; i++) {
u32 u = Eu[i], v = Ev[i];
if (i + 32 < m) {
__builtin_prefetch(&T[gpos[Eu[i + 32] >> LOG2G]], 1, 0);
__builtin_prefetch(&T[gpos[Ev[i + 32] >> LOG2G]], 1, 0);
}
T[gpos[u >> LOG2G]++] = (u32)((u & GMASK) << 20) | (v + 1);
T[gpos[v >> LOG2G]++] = (u32)((v & GMASK) << 20) | (u + 1);
}
}
// per-group scatter into the CSR
if (!fused) for (u32 u = 1; u <= n; u++) SC[u] = start[u];
if (fused) {
// BX2 `shift`: start[] itself is the cursor array (see the PS macro above).
for (u32 g = 0; g < G; g++) {
u32 base = (g << LOG2G) + 1;
u32 e = gpos[g];
// CJ1: the `t + 32 < e` compare+branch is paid on every one of the 2m entries of the
// row's hottest scatter loop, and it protects only the last 32 of each group.
// Same guard split this artifact's sibling row 1006 already uses (H1 in induceSAs0);
// on 1006 the same edit boarded -0.416 %.
{ u32 t = (g + 1) * P;
for (; t + 32 < e; t++) {
u32 x = adj[t];
u32 u = base + (x >> 20);
__builtin_prefetch(&adj[start[base + (adj[t + 32] >> 20)]], 1, 0);
adj[start[u]++] = x & 0xFFFFF;
}
for (; t < e; t++) {
u32 x = adj[t];
u32 u = base + (x >> 20);
adj[start[u]++] = x & 0xFFFFF;
} }
}
} else {
for (u32 g = 0; g < G; g++) {
u32 base = (g << LOG2G) + 1;
u32 hi = base + GSIZE;
if (hi > n + 1) hi = n + 1;
u32 lo = start[base], e = start[hi];
{ u32 t = lo;
for (; t + 32 < e; t++) {
u32 x = T[t];
u32 u = base + (x >> 20);
__builtin_prefetch(&adj[SC[base + (T[t + 32] >> 20)]], 1, 0);
adj[SC[u]++] = x & 0xFFFFF;
}
for (; t < e; t++) {
u32 x = T[t];
u32 u = base + (x >> 20);
adj[SC[u]++] = x & 0xFFFFF;
} }
}
// fallback only: start[] is pristine here (SC was the cursor), so move it one index
// LEFT so the PS() view the readers use is correct.
for (u32 i = 0; i <= n; i++) start[i] = start[i + 1];
}
// sort each adjacency list (fused: the ring case also seeds DV[] and the leaf queue here,
// so deg[] is read once instead of twice and the separate init pass disappears)
u32 qh = 0, qt = 0;
if (m == n - 1) {
if (g_nosort) {
u32 b = 0; // start[1]=0 and start[u+1]=start[u]+deg[u]
for (u32 u = 1; u <= n; u++) { u32 d = deg[u]; bsort(adj + b, d); b += d; }
}
} else {
// deg[] is ALSO the peel's DV[] (see below). The list base is a RUNNING SUM of deg[]
// (start[1]=0, start[u+1]=start[u]+deg[u]), so start[] is never re-read here: the shipped
// loop streamed a full extra 2 MB of start[] to obtain a value it was already holding.
// deg[] is DEAD from the moment this loop ends (the only later readers of a degree
// are the peel, and the peel destroys it) -- so DV[] is not a second array at all.
// The shipped code copied deg[] into DV[] : a full n-vertex 2 MB write for nothing.
if (g_nosort) {
u32 b = 0;
for (u32 u = 1; u <= n; u++) {
u32 d = deg[u];
if (d < 3u) {
/* [p02f-B] branchless for the two dominant list lengths: d == 1 and d == 2 are
88 % of the testcase's lists, and the bsort dispatch is a data-dependent branch
chain (gcc 9.3: `cmp 3 / je / cmp 2 / je / cmp 4 / ...`) -- measured on the judge
as 180 695 mispredicts per 5e5 lists locally, 18.4 % of the whole program's.
The masked form is bit-identical: d == 2 uses the very same `x < y` predicate as
bsort's len == 2 specialisation, and for d < 2 both words are written back with
the values they already hold. adj[b + 1] is in bounds (b + 1 <= 2m < AMAX). */
/* [G1b] one 64-bit pair load/store + byte-rotated copy + a single conditional move
(see REFERENCES): `change` is true exactly for d == 2 with an inverted pair, the
same predicate the masked form used (`x > y` == `!(x < y)` on unequal words; on
equal words neither form swaps). d == 1 leaves `change` 0 => the original pair is
stored back unchanged. adj[b + 1] stays in bounds (b + 1 <= 2m < AMAX). */
u64 pair=*(u64*)(adj+b), swapped=(pair<<32)|(pair>>32);
u32 tt;
__asm__("cmpq %2,%1; movl %3,%0; adcl $0,%0" : "=&r"(tt) : "r"(pair), "r"(swapped), "r"(d) : "cc");
__asm__("cmpl $3,%k1; cmove %2,%0" : "+r"(pair) : "r"(tt), "r"(swapped) : "cc");
*(u64*)(adj+b)=pair;
} else bsort(adj + b, d);
b += d;
// CJ1: `if (d == 1) qq[qt++] = u;` is a ~50/50 data-dependent branch on every one of
// the n vertices. Store unconditionally to qq[qt] (qt <= n < MAXN+16 so it is in
// bounds) and advance the count instead -- a non-qualifying store is overwritten by
// the next qualifier, and the peel/DFS only ever read qq[qh] for qh < qt.
qq[qt] = u;
__asm__("cmpl $2,%k1; adcl $0,%k0" : "+r"(qt) : "r"(d) : "cc");
}
}
}
// DW1a UNI: ONE DFS body for both m == n-1 and m == n. The tree case IS the ring case
// with cu == cv == 0 -- a cut edge no vertex id can match (ids are >= 1) -- so the tree
// path now inherits the ring path's deg-2 chain collapse, its loop-invariant `cutv`
// hoist and its second-level prefetches. Measured on the BOARD: porting ONLY the
// collapse to the tree path took tc12 (a tree case) from 31.622 to 28.885 ms = -8.65 %.
// A single body also stops tree-path edits from re-laying-out the ring path.
u32 cu = 0u, cv = 0u;
if (m != n - 1) {
// ---- m == n : base ring tree ----
while (qh < qt) {
u32 u = qq[qh++];
// Hide the random start[u] miss: issue it 8 pops early. qq[] is sequential, so this
// read is nearly free, and qq[] is zero-filled BSS past qt so the prefetch address is
// always a valid index into start[].
// two-stage lookahead: the pop is a chain qq -> start[u] -> adj[start[u]] -> deg[v].
// Stage 1 (h2) prefetched start[]; stage 2 loads that start 8 pops early (a real load,
// so its value is available to form the adj[] address) and prefetches adj[] from it.
u32 la = qq[qh + 12];
u32 li = la ? la - 1u : 0u;
__builtin_prefetch(&start[li]);
__builtin_prefetch(&adj[start[li]]);
deg[u] = 0;
u32 c=PS(u),v,dv;
do{v=adj[c++];dv=deg[v];}while(!dv);
u32 nd=dv-1;deg[v]=nd;qq[qt]=v;qt+=(nd==1u);
}
u32 c0, br = 0;
if (deg[1]) { c0 = 1; br = 0; }
else {
u32 sp = 0;
VB_SET(1); stk[sp++] = 1;
c0 = 0;
while (sp && !c0) {
u32 u = stk[--sp];
for (u32 c = PS(u); c < PS(u + 1); c++) {
u32 v = adj[c];
if (VB_GET(v)) continue;
if (deg[v]) { c0 = v; br = u; break; }
VB_SET(v); stk[sp++] = v;
}
}
}
u32 a = 0, bb2 = 0;
for (u32 c = PS(c0); c < PS(c0 + 1); c++) {
u32 v = adj[c];
if (deg[v]) { if (!a) a = v; else bb2 = v; }
}
u32 first = a < bb2 ? a : bb2;
u32 lastv = (a == first) ? bb2 : a; // == vs[k-1], known without walking
u32 v1 = first, vk1 = lastv;
u32 backA0 = vk1;
{
// adjacency is sorted ascending => the first non-cycle x in (v1, vk1) IS the minimum.
u32 best = 0;
u32 lo = PS(c0), hi = PS(c0 + 1);
while (lo < hi) { u32 mid = (lo + hi) >> 1; if (adj[mid] <= v1) lo = mid + 1; else hi = mid; }
for (u32 c = lo; c < PS(c0 + 1); c++) {
u32 x = adj[c];
if (x >= vk1) break;
if (deg[x] || x == br) continue;
best = x; break;
}
if (best) backA0 = best;
}
// ---- LAZY FUSED cycle enumeration + `chosen` analysis --------------------
// The shipped code materialised the WHOLE cycle (vs[0..k-1]) and only THEN ran
// the `chosen` loop, which breaks at the first qualifying b. The loop's two
// tests depend only on the pair (vs[b], vs[b+1]), so enumerate one step at a
// time and stop exactly where the loop stops. vs[] is needed for nothing
// else: vs[0]=c0, vs[1]=first, vs[k-1]=lastv are all known up front.
// EQUIVALENCE IS A PROOF, NOT A MEASUREMENT: identical b's in identical order
// with identical conditions, and the cut edge is the same pair in both cases.
{
u32 glast = 0, lastb = 0, chosen = 0;
u32 prev = c0, cur = first, bcnt = 1;
u32 oa = lastv, ob = c0; // the no-break cut: (vs[k-1], vs[0])
while (1) {
u32 nxt = 0;
for (u32 c = PS(cur); c < PS(cur + 1); c++) {
u32 v = adj[c];
if (deg[v] && v != prev) { nxt = v; break; }
}
if (nxt == c0) break; // cur == vs[k-1] => b would be k-1, out of range
u32 gb = 0;
u32 lo = PS(cur), hi = PS(cur + 1);
while (lo < hi) { u32 mid = (lo + hi) >> 1; if (adj[mid] <= nxt) lo = mid + 1; else hi = mid; }
for (u32 c = lo; c < PS(cur + 1); c++) { u32 x = adj[c]; if (deg[x]) continue; gb = x; break; }
if (!gb) {
u32 bbx = lastb ? glast : backA0;
if (bbx < nxt) { oa = cur; ob = nxt; chosen = 1; break; }
} else { glast = gb; lastb = 1; }
prev = cur; cur = nxt; bcnt++;
}
(void)bcnt; (void)chosen;
cu = oa; cv = ob;
}
} // end `if (m != n - 1)`
// ---- the ONE DFS (tree and ring share it) ----
u32 sp = 0;
ofp = op_;
// BW3-a: par[] carries the frame's DFS parent (it is already allocated and the peel is
// finished with it). The parent CANNOT be packed into stk[]: ids need 19 bits and a
// u32 holds only 19+13, so `v | (parent<<19)` overflows for parent > 8191 -- measured as
// a runaway walk that overflows the output buffer (ASan: SEGV in oflush).
stkp[0] = 0; stk[sp] = 1u; stkc[sp] = PS(1); stke[sp] = PS(2); sp++;
emit(1);
while (sp) {
u32 u = stk[sp - 1];
u32 pu = stkp[sp - 1];
u32 c = stkc[sp - 1];
u32 e = stke[sp - 1];
u32 cutv = (u == cu) ? cv : ((u == cv) ? cu : 0xFFFFFFFFu);
// CJ1: guard split (see the fused-path skip loop). If the guarded loop breaks on a
// non-skippable v, the tail re-tests the SAME c and breaks identically -- idempotent.
// DW1a SD: `s` carries PS(adj[c]). The shipped loop PREFETCHES PS(adj[c+1]) -- an
// address, not a value -- so the descend's `PS(v)` and the chain's first `adj[sv]` are
// still cold when we finally take an entry. Loading the list-start one entry early
// does two things at once: `sv` is already in a register when we descend, and
// `prefetch(&adj[s])` puts the list itself in flight a full iteration before it is
// walked. The plain prefetches are kept so the start[] line is still warmed.
u32 s = 0;
if (c < e) s = PS(adj[c]);
while (c + 2 < e) {
u32 v = adj[c];
u32 s1 = PS(adj[c + 1]);
__builtin_prefetch(&adj[s]);
__builtin_prefetch(&adj[s1]);
__builtin_prefetch(&PS(adj[c + 2]));
if (v == pu || v == cutv) { c++; s = s1; }
else break;
}
while (c < e) {
u32 v = adj[c];
__builtin_prefetch(&adj[s]);
if (v == pu || v == cutv) { c++; if (c < e) s = PS(adj[c]); }
else break;
}
if (c < e) {
u32 v = adj[c];
u32 sv = s, sv1 = PS(v + 1);
stkc[sp - 1] = c + 1;
emit(v);
// BW3-a DEG-2 CHAIN COLLAPSE. A deg-2 vertex's two neighbours are its parent u and
// exactly one child w, so the descent is deterministic and the 3 stack stores + the
// pop + a whole outer iteration are pure bookkeeping. u is always one of the two
// because (u,v) is a tree edge -- the cut edge was skipped by the test above.
while (sv1 - sv == 2u) {
u32 q = adj[sv], r = adj[sv + 1];
u32 w = (q == u) ? r : q;
// The original hid this load: each pushed frame's tail did `adj[start[v]]` then
// prefetched start[] for the NEXT iteration. The collapse removes the push, so the
// prefetch must move INTO the chain or the 40-cycle random load is exposed.
/* [p02g-3] the child's list start is loaded HERE, where the shipped body only issued a
prefetch for it: sv/sv1 are consumed only at the bottom of the body, so the random
L3 latency overlaps the cut-edge test and emit(w) instead of serialising with the
following adj[] load. On a break, sv/sv1 keep the current vertex's values. */
u32 nsv = PS(w), nsv1 = PS(w + 1);
__builtin_prefetch(&adj[nsv]);
if ((v == cu && w == cv) || (v == cv && w == cu)) break; /* cut edge: no descent */
emit(w);
u = v; v = w; sv = nsv; sv1 = nsv1;
}
if (sv1 - sv > 1) {
u32 w0 = adj[sv]; __builtin_prefetch(&PS(w0)); __builtin_prefetch(&adj[PS(w0)]);
stkp[sp] = u; stk[sp] = v; stkc[sp] = sv; stke[sp] = sv1; sp++;
}
} else sp--;
}
oflushall();
return;
}
#ifdef LOCAL_TEST
int main() {
size_t sn = fread(g_in, 1, sizeof(g_in), stdin);
ip_ = (const unsigned char *)g_in; ie_ = ip_ + sn;
op_ = g_out;
main_solve();
fwrite(g_out, 1, (size_t)(op_ - g_out), stdout);
return 0;
}
#else
int main() { return 0; }
extern "C" void __libc_start_main(void *mm, int argc, char **argv) {
(void)mm;
unsigned long *p = (unsigned long *)(argv + argc + 1);
while (*p) p++;
p++;
for (; p[0]; p += 2) if (p[0] == 0x6b637564UL) { D = (DI *)p[1]; break; }
if (D) {
ip_ = (const unsigned char *)D->s; ie_ = ip_ + D->sn; op_ = D->o;
main_solve();
D->os = (u64)(op_ - D->o);
}
__asm__ volatile("syscall" ::"a"(60), "D"(0) : "rcx", "r11", "memory");
for (;;);
}
#endif