/* References: saffah_codex_6a_agg3 https://duck.ac/submission/121037,
reused the unique-live-neighbor graph engine and all prior attributions below.
No independent license declared.
Idea: Store each DFS frame's vertex, parent and two CSR cursors together in
one aligned 16-byte record. Four separate stacks accessed the same depth in
four distant cache lines; the new record also permits one vector push store.
Purpose: Reduce DFS stack cache traffic and separate base-address bookkeeping.
*/
/* References: saffah_codex_6a_agg3 https://duck.ac/submission/120868,
reused the graph engine and all inherited attributions below; the external
https://duck.ac/submission/120582 supplies unique-live-neighbor early stopping.
No independent license declared.
Idea: Strengthen the invariant for a connected unicyclic graph: every queued
leaf has exactly one live neighbor when popped, since its remaining component
contains the cycle. Scan until that neighbor is found, eliminating the second
CSR offset load and the redundant adjacency-bound test on each iteration.
Purpose: Reduce leaf-peeling loads and branches while preserving the live core.
*/
/* References: saffah_codex_6a_agg3 https://duck.ac/submission/120676,
three-digit formatting and the fully attributed graph engine below.
No independent license declared; all earlier credits retained.
Idea: Include each three-digit group's leading-zero count in its table word.
The final shift no longer waits for concatenation and a trailing-zero scan.
Purpose: Shorten decimal formatting dependencies while preserving exact bytes.
*/
/* References: saffah_cc_v41_agg1 https://duck.ac/submission/120582:
reused the graph engine, unique-active-neighbor peel and output staging.
saffah_codex_6a_agg3 https://duck.ac/submission/120181: paired decimal input.
All earlier credits remain below; no independent license declared.
Idea: Format six decimal digits from two three-digit lookups in a4KiB table.
One exact reciprocal multiplication splits at1000, replacing the two serial
division reductions and three two-digit table loads.
Purpose: Reduce formatting dependencies in the DFS output loop.
*/
// ===== 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>
#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];
struct alignas(16) DFSFrame{u32 u,p,c,e;};
static DFSFrame frames[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};
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 la=(unsigned)__builtin_ctz(mask), end=(unsigned)__builtin_ctz(mask&(mask-1));
unsigned lb=end-la-1;
__m128i x=_mm_shuffle_epi8(_mm_and_si128(v,_mm_set1_epi8(15)),_mm_load_si128((const __m128i *)(PSH+(la*8+lb)*16)));
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];
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; qt += (d == 1u);
}
}
}
// 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 + 8];
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).
frames[sp++]={1u,0u,PS(1),PS(2)};
emit(1);
while (sp) {
u32 u = frames[sp-1].u;
u32 pu = frames[sp-1].p;
u32 c = frames[sp-1].c;
u32 e = frames[sp-1].e;
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);
frames[sp-1].c=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.
__builtin_prefetch(&PS(w));
if ((v == cu && w == cv) || (v == cv && w == cu)) break; /* cut edge: no descent */
emit(w);
u = v; v = w; sv = PS(v); sv1 = PS(v + 1);
__builtin_prefetch(&adj[sv]);
}
if (sv1 - sv > 1) {
u32 w0 = adj[sv]; __builtin_prefetch(&PS(w0)); __builtin_prefetch(&adj[PS(w0)]);
frames[sp++]={v,u,sv,sv1};
}
} 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