代码

/*
References:
[1] saffah_cc_v41_agg1, https://duck.ac/submission/111412.
Reused the public optimized CSR graph solver, adjacency sorting, cycle handling
and staged output. Its extensive previous attribution comments remain below.
No separate source license is declared in the public submission.
Idea: Decode each two-integer edge line with one SIMD delimiter mask and
length-pair byte shuffle. Packed multiply-adds produce both node numbers
simultaneously, replacing two serial character loops. The last line retains
the scalar reader to avoid a wide load past the input.
Purpose: Verify output equivalence and measure reduced parsing cost online.
*/
// ===== 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], 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};
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;
}
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 * 429497u) >> 32);   // == v / 10000（v <= 500001 全量核证）
  u32 b = v - a * 10000u;
  u32 c = (u32)((b * 5243u) >> 19);   // == b / 100（b <= 9999 全量核证；32 位乘法可直接用 imm）
  u32 e = b - c * 100u;
  u64 w = TA[a] | TC[c] | TE[e];   /* [x02y-E] 三次预移位查表 + 两次 or（原：3 查表 + 2 shl + 3 or + movabs） */
  // 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 = (u32)(__builtin_ctzll(w ^ 0x3030303030303030ULL) >> 3);
  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 e2 = PS(u + 1);
    for (u32 c = PS(u); c < e2; c++) {
      u32 v = adj[c];
      __builtin_prefetch(&deg[adj[c + 1]]);
      // CJ1: two dependent branches per scanned edge of the whole peel.  deg[v] == 0 leaves
      // deg[v] at 0 under BOTH forms (`--deg[v]` is short-circuited), so only the store is
      // made unconditional.
      u32 dv = deg[v];
      u32 nd = dv ? dv - 1u : 0u;
      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.
        __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)]);
        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

评测结果