diff --git a/CLAUDE.md b/CLAUDE.md
index 1be1a475..30d0e675 100644
--- a/CLAUDE.md
+++ b/CLAUDE.md
@@ -68,7 +68,7 @@ implementation of the same op on that workload.
 | Device time | Compare `device_busy_ms`, never wall-clock span |
 | Two questions | `Ratio`: is another kernel faster. `SOL`: how much faster the hardware allows anyone to go, its binding resource (`mem`/`comp`/`lat`) in a `Bound` column. Import the SOL arithmetic and thresholds from the checkout's roofline tool (M5); never re-derive them here |
 | Which page an op lands on | The manifest entry's `family:`, through `_MANIFEST_FAMILY` — an op TileOPs adds needs no change here. One the manifest does not declare falls back to its package, then to keywords |
-| Page order | `DATA_PAGES` follows the API nav's order over the same families: one page per family except `Conv & Pool` (two) and `Other` (Top-k, FFT, mHC, Engram, the rest). `_BENCH_ORDER` in `hooks.py` repeats it — change one, change the other |
+| Page order | `DATA_PAGES`: Elementwise, RoPE, Reduction, Normalization, Conv & Pool, GEMM, Quantization & Dequantization, Attention, MoE, Sampling, Linear Attention, SSM, Other. One page per family except `Conv & Pool` (two) and `Other` (FFT, mHC, Engram, the rest). `TopkSelectorFwdOp` declares `family: attention`, so its row is on Attention, while the API Reference documents it on the Sampling page. The API Reference nav follows it, with FFT, mHC and Engram after SSM and Top-k on the Sampling page. `_BENCH_ORDER` in `hooks.py` repeats it — change all three together |
 | Op order within a page | The order `docs/api/` names them, read by `api_op_order()`. An op no API page names comes last, ranked by verdict |
 | Rows follow the manifest | One row group per manifest label, one row per dtype under it in a `dtype` column. Labels keep the snapshot's order, which is the manifest's; the key above the table repeats it. A row no manifest describes takes its id, trailing dtype names split off, as its label |
 | Workload shapes | The snapshot names a workload but carries no shapes. `scripts/workload_shape.py` reads them from the spec manifest at the commit the benchmark ran on, joined by the `<label>-<dtype>` the benchmark id is built from. A workload the manifest does not declare keeps its id and gets no shapes — never a guessed one |
@@ -78,14 +78,16 @@ implementation of the same op on that workload.
 English at the site root, Chinese under `/zh/`. A Chinese page is a
 `<name>.zh.md` beside the English `<name>.md`, full prose, never an
 `include-markdown` shell. `backends.md`, `torch-compile.md`, everything under
-`performance-guides/memory-bound/` and the two guides under `user-guide/manifest/`
-and `user-guide/dispatch/` were authored in Chinese: edit the `.zh.md` first, then
-bring the English page in line. Everything else goes the other way.
+`performance-guides/memory-bound/`, `blog/`, `user-guide/development.md` and the
+two guides under `user-guide/manifest/` and `user-guide/dispatch/` were authored
+in Chinese: edit the `.zh.md` first, then bring the English page in line.
+Everything else goes the other way.
 
 | Rule | Detail |
 |------|--------|
 | Coverage | Whichever pages have a `.zh.md` — `ls docs/**/*.zh.md` |
 | Never translate | `api/` and `benchmarks/`, both generated; `design/`, mirrored English |
+| Mirrored, translated | `performance-guides/trace-timeline.md` mirrors TileOPs `docs/perf/trace-timeline.md`; its `.zh.md` is a full translation. When upstream changes that file, update the translation in step |
 | Missing translation | Falls back to English at the same URL, and `hooks.py` prepends a "本页暂无中文版" notice. The fallback runs zh → en only: a page that exists only as `.zh.md` leaves its `nav` entry on a missing file and the English sidebar renders a dead link |
 | Figures | A figure with text needs one SVG per language: `img/<name>.svg` for English and `img/<name>.zh.svg` beside it, which the `zh` build picks up for the same reference. Translate the `<text>` nodes and the `aria-label`, keep the geometry. English runs longer than Chinese — grow the `viewBox` rather than let text overflow. The user-guide figures are drawn from sources under `figures/user-guide/` (`<name>.zh.puml`, `<name>.en.puml`, `manifest/overview.py`): edit the source, then run `figures/user-guide/render.sh` |
 | Nav labels | `nav_translations` in the `i18n` plugin block; keep an entry for every `nav` title |
@@ -99,23 +101,41 @@ bring the English page in line. Everything else goes the other way.
 
 ## Nav
 
-`nav` in `mkdocs.yml` is the page list; its six sections run in the reader's
-order, Design last as contributor-facing.
+`nav` in `mkdocs.yml` is the page list: Home, then Blog, then the remaining
+sections in the reader's order, Design last as contributor-facing.
 
 - Add a new page to `nav`, and its label to `nav_translations`.
-- Put a user-facing topic under User Guide.
+- Put a user-facing topic under User Guide, in the group its index lists it in,
+  and keep the nav in the index's order. The index is a plain list per group;
+  `hooks.py` marks it and `extra.css` draws each item as a card.
 - Keep a label short enough for one line in the sidebar; the page's H1 carries
   the full title.
+- Merging a page into another: delete it, and add its old URL to `_REDIRECTS`
+  in `hooks.py` so published links redirect.
 - List a section's `<dir>/index.md` as a bare path with no title. Given a title
   it is promoted anyway and its sidebar row disappears.
 - Leave `toc.integrate` off: the page TOC renders in the right column, and it is
   incompatible with `navigation.indexes`.
 
+## Blog
+
+`docs/blog/`: plain pages, not Material's `blog` plugin. Why: the plugin
+renders no post under `mkdocs-static-i18n` and warns on its archive pages.
+
+- One post is `blog/<slug>.zh.md` and `blog/<slug>.md`, listed in `nav` under
+  Blog and as one link on `blog/index.md`, newest first.
+- A post opens with a short H1 and a one-line subtitle paragraph. `hooks.py`
+  gives that paragraph the `post-subtitle` class that `extra.css` styles; keep
+  classes and attribute lists out of the post's Markdown.
+- No date line, no in-page TOC: the right column carries the TOC.
+
 ## Conventions
 
 - Measure every number and state its conditions. Say when a count will drift.
 - Admonitions (`!!! note`, `!!! warning`) for callouts; relative Markdown links
   for internal cross-references.
+- Write links as plain Markdown. `extra.css` gives every link in running text an
+  arrow, east within the site and north-east off it, and a teal wash on hover.
 - Link to the TileOPs repo rather than duplicating it. A page authored here that
   mirrors upstream content will drift.
 - Gitignored: `site/`, `__pycache__/`, `.cache/`, `TileOPs/`, and
diff --git a/docs/api/index.md b/docs/api/index.md
index ad23c654..acff5727 100644
--- a/docs/api/index.md
+++ b/docs/api/index.md
@@ -13,30 +13,29 @@ op = GemmFwdOp()                        # construct once, reuse
 d = op(a, b)                         # the specialized kernel is built on first call
 ```
 
-The pages are ordered by how much an op composes: the pointwise transforms first, then
-the axis reductions and the normalizations built on them, then the matmul and the
-expert routing over it, then the windowed and spectral transforms, then the
-sequence-model kernels built on all of the above. It is the order `tileops` declares
-its op families in. The exception is Top-k, whose one op is exported from
-`tileops.attention`.
+The pages run in the order a reader meets the ops in a model: the pointwise
+transforms and the positional rotation, then the axis reductions and the
+normalizations built on them, the windowed kernels, the matmul and the
+quantization around it, then attention, the expert routing after it and sampling,
+then the sequence-model kernels. FFT, mHC and Engram follow, and Trace,
+a tool rather than an op, comes last. The Benchmarks pages use the same order.
 
 | Page | What it covers |
 | --- | --- |
 | [Elementwise](elementwise.md) | unary and binary maps, activations, dropout, and the in-place forms |
+| [RoPE](rope.md) | rotary position embedding — NeoX and interleaved layouts, Llama 3.1, YaRN, LongRoPE |
 | [Reduction](reduction.md) | sums, extrema, arg-reductions, cumulative scans, softmax |
 | [Normalization](normalization.md) | RMSNorm, LayerNorm, GroupNorm, BatchNorm and the fused variants |
-| [Quantization](quantization.md) | INT8, FP8 and INT4 quantization, and INT8 dequantization |
-| [Top-k](topk.md) | top-k selection |
-| [GEMM](linear-algebra.md) | dense matmul — plain, batched, and the fp8 variants |
 | [Pooling](pool.md) | average, max and adaptive pooling, with and without indices, plus the chunked sequence mean |
 | [Convolution](convolution.md) | forward convolution over 1D, 2D and 3D inputs |
-| [FFT](fft.md) | the discrete transform |
-| [MoE](moe.md) | the routed mixture-of-experts FFN and its separately callable stages |
-| [Sampling](sampling.md) | logits masks (top-k, top-p, min-p) and sampling, including chain speculative sampling |
-| [RoPE](rope.md) | rotary position embedding — NeoX and interleaved layouts, Llama 3.1, YaRN, LongRoPE |
+| [GEMM](linear-algebra.md) | dense matmul — plain, batched, and the fp8 variants |
+| [Quantization & Dequantization](quantization.md) | INT8, FP8 and INT4 quantization, and INT8 dequantization |
 | [Attention](attention.md) | forward and backward attention, including the paged and decode kernels |
+| [MoE](moe.md) | the routed mixture-of-experts FFN and its separately callable stages |
+| [Top-k & Sampling](sampling.md) | top-k selection, logits masks (top-k, top-p, min-p) and sampling, including chain speculative sampling |
 | [Linear Attention](linear-attention.md) | DeltaNet, Gated DeltaNet and gated linear attention |
 | [Mamba](mamba.md) | the SSD scan, its decode step, and the chunked forms |
+| [FFT](fft.md) | the discrete transform |
 | [mHC](mhc.md) | Manifold-Constrained Hyper-Connections — the pre/post pair around a layer |
 | [Engram](engram.md) | the Engram GateConv pair and its decode step |
 | [Trace](trace.md) | the in-kernel timeline tracer, a tool rather than an op |
diff --git a/docs/api/quantization.md b/docs/api/quantization.md
index abc25550..5a869add 100644
--- a/docs/api/quantization.md
+++ b/docs/api/quantization.md
@@ -1,4 +1,4 @@
-# Quantization Operators
+# Quantization and Dequantization Operators
 
 Every op on this page is used the same way: construct it once, then call it. The
 constructor takes what the kernel is compiled with; the call takes the tensors.
diff --git a/docs/api/sampling.md b/docs/api/sampling.md
index ed530d2c..138a2998 100644
--- a/docs/api/sampling.md
+++ b/docs/api/sampling.md
@@ -1,13 +1,22 @@
-# Sampling Operators
+# Top-k and Sampling Operators
 
 Every op on this page is used the same way: construct it once, then call it. The
 constructor takes what the kernel is compiled with; the call takes the tensors.
 Both are documented under each op — `__init__` and `forward`, where `forward` is
 what runs when you call `op(...)`.
 
-The masks set every logit a filter drops to `-inf`, so a softmax over the result
-renormalizes over what is kept. The samplers draw from a distribution, and take the
-random seed and offset as tensors so a draw is reproducible.
+The top-k selector returns the indices of the largest scores in each row. The masks
+set every logit a filter drops to `-inf`, so a softmax over the result renormalizes
+over what is kept. The samplers draw from a distribution, and take the random seed
+and offset as tensors so a draw is reproducible.
+
+## Top-k selection
+
+::: tileops.attention.TopkSelectorFwdOp
+    options:
+      show_root_heading: true
+      heading_level: 3
+      members: ["__init__", "forward"]
 
 ## Logits masks
 
diff --git a/docs/api/topk.md b/docs/api/topk.md
deleted file mode 100644
index 82207e5c..00000000
--- a/docs/api/topk.md
+++ /dev/null
@@ -1,14 +0,0 @@
-# Top-k Operators
-
-Every op on this page is used the same way: construct it once, then call it. The
-constructor takes what the kernel is compiled with; the call takes the tensors.
-Both are documented under each op — `__init__` and `forward`, where `forward` is
-what runs when you call `op(...)`.
-
-## Top-k selection
-
-::: tileops.attention.TopkSelectorFwdOp
-    options:
-      show_root_heading: true
-      heading_level: 3
-      members: ["__init__", "forward"]
diff --git a/docs/assets/extra.css b/docs/assets/extra.css
index b323dcd1..e358e620 100644
--- a/docs/assets/extra.css
+++ b/docs/assets/extra.css
@@ -438,6 +438,54 @@ body,
   color: var(--tf-teal);
 }
 
+/* Every link in running text carries an arrow, so it reads as a link before it
+ * is hovered: east for a page on this site, north-east for one that leaves it.
+ * An absolute link back to this site counts as one on it. Hover lifts a teal wash from the underline and nudges the arrow. Scoped to
+ * prose — paragraphs, list items, table cells — so headings' permalinks, the
+ * nav, buttons, badges and the page cards keep their own look, and no page's
+ * Markdown has to mark a link for it.
+ */
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a:not(.md-button, :has(img)) {
+  padding: 0 0.08em;
+  margin: 0 -0.08em;
+  border-radius: 0.15em;
+  background: linear-gradient(var(--tf-teal-tint), var(--tf-teal-tint)) 0 100% / 100% 0 no-repeat;
+  transition: background-size 160ms ease, text-decoration-color 160ms ease;
+}
+
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a:not(.md-button, :has(img)):is(:hover, :focus-visible) {
+  background-size: 100% 100%;
+}
+
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a:not(.md-button, :has(img))::after {
+  content: "";
+  display: inline-block;
+  width: 0.62em;
+  height: 0.62em;
+  margin-left: 0.18em;
+  vertical-align: 0.06em;
+  background-color: currentcolor;
+  -webkit-mask: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20viewBox='0%200%2012%2012'%3E%3Cpath%20d='M2%206h8M6.5%202.5%2010%206l-3.5%203.5'%20fill='none'%20stroke='%23000'%20stroke-width='1.6'%20stroke-linecap='round'%20stroke-linejoin='round'/%3E%3C/svg%3E") center / contain no-repeat;
+  mask: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20viewBox='0%200%2012%2012'%3E%3Cpath%20d='M2%206h8M6.5%202.5%2010%206l-3.5%203.5'%20fill='none'%20stroke='%23000'%20stroke-width='1.6'%20stroke-linecap='round'%20stroke-linejoin='round'/%3E%3C/svg%3E") center / contain no-repeat;
+  opacity: 0.7;
+  transition: transform 160ms ease, opacity 160ms ease;
+}
+
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a[href^="http"]:not([href^="https://tile-ai.github.io/TileOPs.github.io/"], .md-button, :has(img))::after {
+  vertical-align: 0.12em;
+  -webkit-mask-image: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20viewBox='0%200%2012%2012'%3E%3Cpath%20d='M4%202h6v6M10%202%202.5%209.5'%20fill='none'%20stroke='%23000'%20stroke-width='1.6'%20stroke-linecap='round'%20stroke-linejoin='round'/%3E%3C/svg%3E");
+  mask-image: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20viewBox='0%200%2012%2012'%3E%3Cpath%20d='M4%202h6v6M10%202%202.5%209.5'%20fill='none'%20stroke='%23000'%20stroke-width='1.6'%20stroke-linecap='round'%20stroke-linejoin='round'/%3E%3C/svg%3E");
+}
+
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a:not(.md-button, :has(img)):is(:hover, :focus-visible)::after {
+  opacity: 1;
+  transform: translateX(0.1em);
+}
+
+.md-typeset :is(p, li, td, dd):not(.page-cards > li) a[href^="http"]:not([href^="https://tile-ai.github.io/TileOPs.github.io/"], .md-button, :has(img)):is(:hover, :focus-visible)::after {
+  transform: translate(0.08em, -0.08em);
+}
+
 /* A rule between sections, drawn as light rather than as a line. */
 .md-typeset hr {
   height: 1px;
@@ -1297,13 +1345,14 @@ html[lang="zh"] .md-nav__title {
   }
 }
 
-/* Labels in one column, what varies in the next. A label past the column's
- * width wraps at its hyphens, so a long one never pushes the values off.
+/* One column: each label on its own line, what it ran on under it. Two columns
+ * put a label and its shapes side by side, and the eye then reads across a gap
+ * that changes width from one key to the next.
  */
 .md-typeset .wl-key ul.wl-rows {
   display: grid;
-  grid-template-columns: fit-content(24em) minmax(0, 1fr);
-  gap: 0.1em 0.9em;
+  grid-template-columns: minmax(0, 1fr);
+  gap: 0;
   align-items: baseline;
   margin: 0;
   padding: 0;
@@ -1314,11 +1363,12 @@ html[lang="zh"] .md-nav__title {
   display: contents;
 }
 
-/* No width for two columns: each label on its own line, what varies under it. */
-@media (max-width: 50rem) {
-  .md-typeset .wl-key ul.wl-rows {
-    grid-template-columns: minmax(0, 1fr);
-  }
+.md-typeset .wl-key ul.wl-rows .wl-id {
+  margin-top: 0.4em;
+}
+
+.md-typeset .wl-key ul.wl-rows li:first-child .wl-id {
+  margin-top: 0;
 }
 
 /* A hairline between the label and what it ran on — the same rule the table
@@ -1341,20 +1391,6 @@ html[lang="zh"] .md-nav__title {
   min-width: 0;
 }
 
-/* Labels too long for the label column: each on its own line, values under it. */
-.md-typeset .wl-key ul.wl-rows.wl-long {
-  grid-template-columns: minmax(0, 1fr);
-  row-gap: 0;
-}
-
-.md-typeset .wl-key ul.wl-rows.wl-long .wl-id {
-  margin-top: 0.4em;
-}
-
-.md-typeset .wl-key ul.wl-rows.wl-long li:first-child .wl-id {
-  margin-top: 0;
-}
-
 .md-typeset .wl-key .wl-delta:empty {
   border-left: 0;
 }
@@ -1670,3 +1706,143 @@ html[lang="zh"] .md-nav__title {
   line-height: 1.7;
   text-align: left;
 }
+
+/* A blog post's subtitle: the line under the H1 that says what the post is
+ * about, so the H1 itself stays one short line.
+ */
+.md-typeset .post-subtitle {
+  margin: -0.6rem 0 1.6rem;
+  color: var(--tf-muted);
+  font-family: var(--tf-display);
+  font-size: 1.15rem;
+  font-weight: 600;
+  line-height: 1.35;
+}
+
+/* A section index's page list, drawn as cards by topic. `hooks.py` marks the
+ * lists; each item is a link, which becomes the card's title, and a line of
+ * description under it. Groups alternate violet and teal: a faint fill of that
+ * colour, a rule along the card's top, a numbered kicker, and a small corner of
+ * tiles that fills in on hover. Two columns where the page
+ * has the width, one below.
+ */
+.md-typeset:has(> ul.page-cards) {
+  counter-reset: page-card;
+}
+
+.md-typeset ul.page-cards {
+  --card-accent: var(--tf-violet);
+  --card-tint: rgba(90, 62, 133, 0.08);
+  display: grid;
+  grid-template-columns: repeat(auto-fill, minmax(17rem, 1fr));
+  gap: 1rem;
+  margin: 0.9rem 0 1.8rem;
+  padding: 0;
+  list-style: none;
+}
+
+.md-typeset ul.page-cards:nth-of-type(even) {
+  --card-accent: var(--tf-teal);
+  --card-tint: rgba(22, 112, 107, 0.08);
+}
+
+.md-typeset ul.page-cards > li {
+  position: relative;
+  overflow: hidden;
+  min-height: 7.2rem;
+  margin: 0;
+  padding: 1.05rem 1.2rem 1.9rem;
+  border: 1px solid var(--tf-line);
+  border-top: 3px solid var(--card-accent);
+  border-radius: 0.6rem;
+  background: var(--card-tint);
+  box-shadow: 0 1px 2px rgba(27, 19, 39, 0.05);
+  color: var(--tf-muted);
+  font-size: 0.74rem;
+  line-height: 1.6;
+  counter-increment: page-card;
+  transition: border-color 180ms ease, box-shadow 180ms ease, transform 180ms ease;
+}
+
+/* The kicker: the card's number in the index, in the group's accent. */
+.md-typeset ul.page-cards > li::before {
+  content: counter(page-card, decimal-leading-zero);
+  display: block;
+  margin-bottom: 0.45rem;
+  color: var(--card-accent);
+  font-family: var(--tf-mono);
+  font-size: 0.62rem;
+  font-weight: 700;
+  letter-spacing: 0.12em;
+}
+
+/* A corner of tiles, faint until the card is hovered. */
+.md-typeset ul.page-cards > li::after {
+  content: "";
+  position: absolute;
+  top: 0.9rem;
+  right: 0.9rem;
+  width: 46px;
+  height: 46px;
+  background-color: var(--card-accent);
+  -webkit-mask: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20width='46'%20height='46'%3E%3Crect%20x='0'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.9'/%3E%3Crect%20x='12'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.5'/%3E%3Crect%20x='0'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.5'/%3E%3Crect%20x='24'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.25'/%3E%3Crect%20x='12'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.3'/%3E%3Crect%20x='0'%20y='24'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.25'/%3E%3Crect%20x='36'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.12'/%3E%3Crect%20x='24'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.15'/%3E%3Crect%20x='12'%20y='24'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.15'/%3E%3Crect%20x='0'%20y='36'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.12'/%3E%3C/svg%3E") right top / contain no-repeat;
+  mask: url("data:image/svg+xml,%3Csvg%20xmlns='http://www.w3.org/2000/svg'%20width='46'%20height='46'%3E%3Crect%20x='0'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.9'/%3E%3Crect%20x='12'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.5'/%3E%3Crect%20x='0'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.5'/%3E%3Crect%20x='24'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.25'/%3E%3Crect%20x='12'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.3'/%3E%3Crect%20x='0'%20y='24'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.25'/%3E%3Crect%20x='36'%20y='0'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.12'/%3E%3Crect%20x='24'%20y='12'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.15'/%3E%3Crect%20x='12'%20y='24'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.15'/%3E%3Crect%20x='0'%20y='36'%20width='10'%20height='10'%20rx='1.5'%20fill='%23000'%20fill-opacity='0.12'/%3E%3C/svg%3E") right top / contain no-repeat;
+  opacity: 0.18;
+  transition: opacity 180ms ease, transform 180ms ease;
+  pointer-events: none;
+}
+
+.md-typeset ul.page-cards > li:hover,
+.md-typeset ul.page-cards > li:focus-within {
+  border-color: var(--card-accent);
+  box-shadow: 0 14px 30px -18px rgba(27, 19, 39, 0.45);
+  transform: translateY(-3px);
+}
+
+.md-typeset ul.page-cards > li:hover::after,
+.md-typeset ul.page-cards > li:focus-within::after {
+  opacity: 0.6;
+  transform: scale(1.06);
+}
+
+/* The title's link covers the whole card, so the card is one click target. */
+.md-typeset ul.page-cards > li > a:first-child {
+  display: block;
+  max-width: calc(100% - 3.2rem);
+  margin-bottom: 0.35rem;
+  color: var(--md-default-fg-color);
+  font-family: var(--tf-display);
+  font-size: 0.95rem;
+  font-weight: 700;
+  line-height: 1.3;
+  text-decoration: none;
+}
+
+.md-typeset ul.page-cards > li > a:first-child::before {
+  content: "";
+  position: absolute;
+  inset: 0;
+  z-index: 1;
+  border-radius: inherit;
+}
+
+/* The arrow sits in the card's bottom-right corner, so a title that wraps never
+ * leaves it alone on a line.
+ */
+.md-typeset ul.page-cards > li > a:first-child::after {
+  content: "→";
+  position: absolute;
+  right: 1.2rem;
+  bottom: 0.9rem;
+  color: var(--card-accent);
+  font-size: 1rem;
+  transition: transform 180ms ease;
+}
+
+.md-typeset ul.page-cards > li:hover > a:first-child::after {
+  transform: translateX(0.25em);
+}
+
+html[lang="zh"] .md-typeset ul.page-cards > li > a:first-child {
+  font-weight: 700;
+}
diff --git a/docs/blog/index.md b/docs/blog/index.md
new file mode 100644
index 00000000..f735cf6c
--- /dev/null
+++ b/docs/blog/index.md
@@ -0,0 +1,3 @@
+# Blog
+
+Technical explorations from building TileOPs.
diff --git a/docs/blog/index.zh.md b/docs/blog/index.zh.md
new file mode 100644
index 00000000..428dadf7
--- /dev/null
+++ b/docs/blog/index.zh.md
@@ -0,0 +1,3 @@
+# 博客
+
+TileOPs 开发中的技术探索。
diff --git a/docs/design/index.md b/docs/design/index.md
index 015864f0..42e8e830 100644
--- a/docs/design/index.md
+++ b/docs/design/index.md
@@ -4,11 +4,30 @@ Architecture and design documentation for TileOPs internals. The pages
 below mirror `docs/design/` in the [`tile-ai/TileOPs`](https://github.com/tile-ai/TileOPs)
 repository — the source of truth — pulled in at site build time.
 
-- [Architecture](architecture.md) — top-level module layout and the spec-driven pipeline.
-- [Op Manifest](manifest.md) — the `src/tileops/manifest/` package as the source of truth for op interfaces.
-- [Op Interface Design](ops-design.md) — playbook for scaffolding a new op from a manifest entry.
-- [Op Interface Reference](ops-design-reference.md) — interface tables, codegen, naming, and the family-base protocol.
-- [Slot Rules](op-slot-rules.md) — the authoritative rule, example, and common mistakes per op-file slot.
-- [Roofline](roofline.md) — performance model and the `roofline` manifest field.
-- [Testing & Benchmarking](testing.md) — separation of correctness tests and profiling benchmarks.
-- [Layer Boundaries](layer-boundaries.md) — what each layer owns and depends on, and the interfaces layers compose through.
+## System structure
+
+- [Architecture](architecture.md)
+  Top-level module layout and the spec-driven pipeline.
+- [Layer Boundaries](layer-boundaries.md)
+  What each layer owns and depends on, and the interfaces layers compose through.
+
+## The spec
+
+- [Op Manifest](manifest.md)
+  The `src/tileops/manifest/` package as the source of truth for op interfaces.
+
+## Building an op
+
+- [Op Interface Design](ops-design.md)
+  Playbook for scaffolding a new op from a manifest entry.
+- [Op Interface Reference](ops-design-reference.md)
+  Interface tables, codegen, naming, and the family-base protocol.
+- [Slot Rules](op-slot-rules.md)
+  The authoritative rule, example, and common mistakes per op-file slot.
+
+## Acceptance
+
+- [Testing & Benchmarking](testing.md)
+  Separation of correctness tests and profiling benchmarks.
+- [Roofline](roofline.md)
+  Performance model and the `roofline` manifest field.
diff --git a/docs/index.md b/docs/index.md
index 57f089a4..74686c85 100644
--- a/docs/index.md
+++ b/docs/index.md
@@ -1,25 +1,23 @@
 # TileOPs
 
-TileOPs is an operator library for large-model inference, built on
-[TileLang](https://github.com/tile-ai/tilelang). One set of op interfaces can be
-implemented by different backends on different hardware.
-
-TileOPs differs from a hand-written operator library in how it is organised: every
-op is first declared as a spec, and an agent then generates the implementation from
-that spec. The spec is the only input to code generation and the standard the result
-is accepted against:
-
-- correctness is judged against the reference implementation the spec names;
-- performance is judged against the bound the roofline model gives.
-
-Neither check depends on human judgement. An implementation can therefore be
-regenerated from its spec at any time, while a spec cannot be derived from an
-implementation.
-
-To a caller, TileOPs is a set of ops that can be called directly. Shapes and dtype
-are fixed at call time; the specialized kernel is built and cached on first use and
-can then be used with CUDA graphs. Each op declares whether it supports
-`torch.compile(fullgraph=True)`.
+TileOPs is an exploratory operator library for large-model inference, built on
+[TileLang](https://github.com/tile-ai/tilelang). TileOPs is designed for agents,
+and agents build the whole project. Such a project has to meet concrete
+code-quality requirements: its structure stays consistent, it does not diverge or
+bloat as ops are added, and its code stays maintainable. The design of TileOPs
+serves three goals:
+
+- **Maintainable.** Each op is declared by a spec, and an agent generates the
+  implementation from it. The ops in a family share one set of interfaces and
+  rules, and every new op and kernel follows them.
+- **Verifiable.** The spec names the reference implementation correctness is
+  judged against. Tests compare a kernel's output with that reference, and
+  performance measurements compare its measured speed with the bound the roofline
+  model gives. The acceptance criteria are fixed in advance, and the checks run
+  automatically.
+- **Tunable.** The roofline model gives the gap between each kernel and its
+  performance bound. The [nightly benchmarks](benchmarks/index.md) compare each
+  kernel with the fastest other implementation on the same hardware.
 
 ## Installation
 
@@ -46,6 +44,7 @@ flops, nbytes = op.eval_roofline()   # what the call had to do and move
 
 ## Where to go next
 
+- [Blog](blog/index.md): technical explorations from building TileOPs.
 - [User Guide](user-guide/index.md): reading and writing the manifest, bringing an op
   into `torch.compile`, how a benchmark is timed, and adding a hardware backend.
 - [API Reference](api/index.md): the constructor parameters and call signatures of
diff --git a/docs/index.zh.md b/docs/index.zh.md
index 1eaf4269..f0fe0df5 100644
--- a/docs/index.zh.md
+++ b/docs/index.zh.md
@@ -1,15 +1,10 @@
 # TileOPs
 
-TileOPs 是一个面向大模型推理的算子库，构建在 [TileLang](https://github.com/tile-ai/tilelang) 之上。同一套 op 接口可以由不同 backend 在不同硬件上实现。
+TileOPs 是一个面向大模型推理的探索性算子库，构建在 [TileLang](https://github.com/tile-ai/tilelang) 之上。TileOPs 为 agent 设计，整个项目由 agent 构建。这样的项目需要满足明确的代码质量要求：项目结构保持一致，不随 op 增多而发散或膨胀，代码可以长期维护。TileOPs 的设计围绕三个目标：
 
-TileOPs 与手写算子库的区别在于组织方式：每个 op 先以一份 spec 声明，再由 agent 依据这份 spec 生成实现。spec 是代码生成的唯一依据，也是验收的标准：
-
-- 正确性以 spec 指定的参考实现为准；
-- 性能以 roofline 模型给出的上界为准。
-
-两项验收都不依赖人的判断。因此一个实现可以随时从 spec 重新生成，spec 却不能从实现反推。
-
-对使用者而言，TileOPs 提供一批可以直接调用的 op。形状与 dtype 在调用时确定；特化后的 kernel 在首次使用时构造并缓存，之后可以与 CUDA graph 配合使用。每个 op 各自声明是否支持 `torch.compile(fullgraph=True)`。
+- **可维护**：每个 op 由一份 spec 声明，agent 依据 spec 生成实现。同一 family 的 op 共用一套接口与规则，新增的 op 和 kernel 遵循这套接口与规则。
+- **可验证**：spec 指定正确性所依据的参考实现。测试将 kernel 的输出与参考实现比较，性能评测将实测性能与 roofline 模型给出的上界比较。验收标准事先确定，检查过程自动执行。
+- **可调优**：roofline 模型给出每个 kernel 与性能上界之间的差距。[每晚的 benchmark](benchmarks/index.md) 在同一硬件上将每个 kernel 与最快的其他实现比较。
 
 ## 安装
 
@@ -35,7 +30,8 @@ flops, nbytes = op.eval_roofline()   # 本次调用所需的计算量与访存
 
 ## 后续阅读
 
-- [使用指南](user-guide/index.md)：读写 manifest、接入 `torch.compile`、benchmark 的计时方法、接入新硬件 backend。
+- [博客](blog/index.md)：TileOPs 开发中的技术探索。
+- [用户指南](user-guide/index.md)：读写 manifest、接入 `torch.compile`、benchmark 的计时方法、接入新硬件 backend。
 - [API 参考](api/index.md)：各 op family 的构造参数与调用方式。
 - [性能数据](benchmarks/index.md)：每晚在 H200 上实测，逐个 workload 与其他实现对比。
 
diff --git a/docs/performance-guides/memory-bound/global-memory-access.md b/docs/performance-guides/memory-bound/global-memory-access.md
index 45631d6b..8966d00d 100644
--- a/docs/performance-guides/memory-bound/global-memory-access.md
+++ b/docs/performance-guides/memory-bound/global-memory-access.md
@@ -6,8 +6,8 @@ choosing one.
 
 ## Checking whether DRAM bandwidth is the current limit {#regime}
 
-[Elementwise](https://tile-ai.github.io/TileOPs.github.io/api/elementwise/) and
-[Reduction](https://tile-ai.github.io/TileOPs.github.io/api/reduction/) are the
+[Elementwise](../../api/elementwise.md) and
+[Reduction](../../api/reduction.md) are the
 typical memory-bound kernels. Each recommendation on this page states when it
 applies, why it applies, and what the wrong and right code look like.
 
diff --git a/docs/performance-guides/memory-bound/global-memory-access.zh.md b/docs/performance-guides/memory-bound/global-memory-access.zh.md
index f10cc19d..4893d8e3 100644
--- a/docs/performance-guides/memory-bound/global-memory-access.zh.md
+++ b/docs/performance-guides/memory-bound/global-memory-access.zh.md
@@ -4,7 +4,7 @@
 
 ## 确认 DRAM 带宽是否为当前的瓶颈 {#regime}
 
-[Elementwise](https://tile-ai.github.io/TileOPs.github.io/api/elementwise/) 与 [Reduction](https://tile-ai.github.io/TileOPs.github.io/api/reduction/) 是典型的访存受限 kernel。本页每条建议都写明触发条件、成因，以及反例与正例代码。
+[Elementwise](../../api/elementwise.md) 与 [Reduction](../../api/reduction.md) 是典型的访存受限 kernel。本页每条建议都写明触发条件、成因，以及反例与正例代码。
 
 本页的实测都在同一组条件下取得：**输入大于 L2 的 60 MiB，且 block 数足以填满整卡**（H200 有 132 个 SM）。此时 DRAM 带宽是主要瓶颈，访存模式的差别直接反映在性能上。
 
diff --git a/docs/performance-guides/trace-timeline.zh.md b/docs/performance-guides/trace-timeline.zh.md
new file mode 100644
index 00000000..1958dd11
--- /dev/null
+++ b/docs/performance-guides/trace-timeline.zh.md
@@ -0,0 +1,121 @@
+# kernel 内的时间线追踪 { #in-kernel-timeline-trace }
+
+[`tileops.trace`](../api/trace.md) 是一个在 kernel 内部记录时间线的追踪工具，用于诊断 kernel 的性能。它在 kernel 内为每个 CTA 记录时间戳，并渲染成一条可以滚动查看的时间线，显示执行区间、空隙，以及 producer 与 consumer 之间的重叠。这些信息是 `ncu` 这类以 kernel 为单位的 profiler 看不到的。warp specialization 的 kernel 最能用上它：这类 kernel 的关键就是让 producer（TMA）与 consumer（WGMMA）两个 warpgroup 的执行相互重叠。
+
+## 工作方式 { #how-it-works }
+
+1. 在 kernel 代码中用标记（`trace.range`、`trace.group` 等）标注要记录的区间。
+1. 标记**总是以占位的形式生成**。构建时，kernel 要么被 **lower**：标记变成真正调用 `clock64()` 记录时间的代码，并在输出末尾多出一个 `slots`；要么被 **strip**：标记变成空操作，生成的 CUDA 与未插桩的版本完全相同。
+1. 运行时，`trace.run` 执行 kernel，解码 `slots` 缓冲区，并写出一个自包含的 Plotly HTML 时间线。
+1. 由进程内的开关 `trace.enable()` 决定 lower 还是 strip，因此关闭时追踪**没有任何开销**，标记可以留在生产代码中。
+
+时间戳来自 `clock64()`，即每个 SM 的周期计数器。
+
+## 编写带追踪的 kernel { #write-a-traced-kernel }
+
+下面是一个完整的、接入了追踪的 warp specialization GEMM（示意用，单缓冲）。编号 `(1)`–`(7)` 的标记是仅有的与追踪相关的代码，各自的说明和 API 文档链接见代码下方。生产环境中的多 stage 版本见 [`src/tileops/kernels/gemm/dense.py`](https://github.com/tile-ai/TileOPs/blob/main/src/tileops/kernels/gemm/dense.py)。
+
+```{ .python .annotate }
+import functools
+import tilelang
+import tilelang.language as T
+from tileops.trace import trace  # (1)!
+
+
+@functools.lru_cache(maxsize=32)
+def build_gemm(m, n, k, dtype="float16", traced=False):
+    @tilelang.jit(out_idx=trace.out_idx(1, traced))  # (2)!
+    def factory(block_m=128, block_n=128, block_k=64):
+        @T.prim_func
+        def main(a: T.Tensor((m, k), dtype), b: T.Tensor((n, k), dtype),
+                 c: T.Tensor((m, n), dtype)):
+            with T.Kernel(T.ceildiv(n, block_n), T.ceildiv(m, block_m),
+                          threads=256) as (bx, by):
+                a_smem = T.alloc_shared((block_m, block_k), dtype)
+                b_smem = T.alloc_shared((block_n, block_k), dtype)
+                c_local = T.alloc_fragment((block_m, block_n), "float")
+                full = T.alloc_barrier(128)
+                tx = T.get_thread_binding()
+
+                if tx < 128:
+                    with trace.group("producer", lead=0):  # (3)!
+                        for ki in T.serial(T.ceildiv(k, block_k)):
+                            with trace.range("tma", lane="tma"):  # (4)!
+                                T.tma_copy(a[by * block_m, ki * block_k], a_smem, barrier=full)
+                                T.tma_copy(b[bx * block_n, ki * block_k], b_smem, barrier=full)
+                            with trace.range("arrive", lane="barrier"):
+                                T.barrier_arrive(full)
+                else:
+                    with trace.group("consumer", lead=128):
+                        T.clear(c_local)
+                        for ki in T.serial(T.ceildiv(k, block_k)):
+                            with trace.range("wait", lane="barrier"):
+                                T.barrier_wait(full, ki % 2)
+                            with trace.range("mma", lane="wgmma"):  # (5)!
+                                T.wgmma_gemm(a_smem, b_smem, c_local, transpose_B=True)
+                        with trace.range("epilogue"):
+                            T.copy(c_local, c[by * block_m, bx * block_n])
+
+                trace.dag("arrive", "wait")  # (6)!
+        return trace.finalize(main, traced=traced, max_events=1024)  # (7)!
+    return factory
+```
+
+1. 导入 trace 命名空间。下面的每个调用都是这个 `trace` 对象上的方法，完整说明见 [API 参考](../api/trace.md)。
+1. [`trace.out_idx(n_outputs, traced)`](../api/trace.md#tileops.trace.api._Trace.out_idx) 给出 `@tilelang.jit` 的 `out_idx`。只有 `traced` 时它才多出一个位置给末尾的 `slots` 输出，因此同一个构建函数在开启和关闭追踪时都能使用。
+1. [`trace.group(name, lead)`](../api/trace.md#tileops.trace.api._Trace.group) 声明由哪个 warpgroup 记录。`lead` 是被选出的写入线程（`tx == lead`）：计算仍在所有线程上执行，只有时间戳由 `lead` 写入。
+1. [`trace.range(name, lane)`](../api/trace.md#tileops.trace.api._Trace.range) 是一个 `with` 块，从进入计时到退出，在子行 `lane` 上画成一段条形。控制流放不进 `with` 时，用 [`trace.range_start`](../api/trace.md#tileops.trace.api._Trace.range_start) 与 [`trace.range_end`](../api/trace.md#tileops.trace.api._Trace.range_end)；只需要一个零宽度的时间点时，用 [`trace.record`](../api/trace.md#tileops.trace.api._Trace.record)。
+1. lane 的名字（`"tma"`、`"barrier"`、`"wgmma"`，以及默认的 `"main"`）就是时间线上的各行。
+1. [`trace.dag(src, dst)`](../api/trace.md#tileops.trace.api._Trace.dag) 声明从一个命名区间指向另一个区间的依赖箭头（`arrive` → `wait`），每出现一次画一条。
+1. [`trace.finalize(func, traced, max_events)`](../api/trace.md#tileops.trace.api._Trace.finalize) 在 `traced` 时 lower 标记并加上 `slots` 输出，否则把标记 strip 成零开销。`traced` **必须是构建函数缓存键的一部分**，这样同一形状下带追踪和不带追踪的两次构建才不会冲突。
+
+## 运行带追踪的 kernel { #running-a-traced-kernel }
+
+调用构建函数时传入 `traced=trace.enabled`，再把编译好的 kernel 交给 `trace.run`。同一个 `forward` 在两种模式下都能使用：追踪关闭时原样返回输出；追踪开启时写出时间线，并且只返回真正的输出。因此调用方不需要自己按开关分支。
+
+```{ .python .annotate }
+def forward(self, a, b):
+    compiled = build_gemm(self.m, self.n, self.k, self.dtype_str,
+                          traced=trace.enabled)(**self.config)  # (1)!
+    return trace.run(compiled, (a, b), stem="gemm_128x256x512")  # (2)!
+```
+
+1. 按开关构建对应的 kernel：[`trace.enabled`](../api/trace.md#tileops.trace.api._Trace.enabled) 决定选用带追踪的版本还是 strip 后的版本，它也是标记 `(7)` 所说的缓存键的一部分。
+1. [`trace.run(compiled, inputs, stem=...)`](../api/trace.md#tileops.trace.api._Trace.run) 运行 kernel。带追踪时，它把末尾的 `slots` 拆出来解码，写出 `debug/<stem>.html`，每次调用都生成一个不会重名的新文件。它内部由 [`decode`](../api/trace.md#tileops.trace.api._Trace.decode) 和 [`dump`](../api/trace.md#tileops.trace.api._Trace.dump) 组成，两者也可以直接调用。
+
+## 开启追踪 { #enabling-tracing }
+
+追踪默认关闭。在程序启动时打开一次进程内的开关，之后照常运行：
+
+```{ .python .annotate }
+from tileops.trace import trace
+
+trace.enable()  # (1)!
+c = op.forward(a, b)  # (2)!
+```
+
+1. [`trace.enable(output="debug")`](../api/trace.md#tileops.trace.api._Trace.enable) 打开追踪，并指定输出目录，默认是 `debug/`（已被 gitignore）。这个开关只在当前进程内有效：不读取环境变量，也不对 `tilelang` 做 monkeypatch。相关的还有 [`trace.disable()`](../api/trace.md#tileops.trace.api._Trace.disable)、[`trace.enabled`](../api/trace.md#tileops.trace.api._Trace.enabled) 与 [`trace.output`](../api/trace.md#tileops.trace.api._Trace.output)。
+1. 此后运行的每个带追踪的 kernel 都会写出 `debug/<stem>.html`。
+
+在 pytest 中，`--trace-kernel` 会在任何 kernel 构建之前，由 `pytest_configure` 调用 `trace.enable()`：
+
+```bash
+pytest tests/ops/test_gemm.py --trace-kernel
+```
+
+## 阅读时间线 { #reading-the-timeline }
+
+<iframe src="../../../performance-guides/gemm-trace.html" title="GEMM 时间线" width="100%" height="540"
+        style="border:1px solid var(--md-default-fg-color--lightest);border-radius:4px;"></iframe>
+
+- **顶部的 CTA 标签页**：每个 CTA（block）一条时间线。
+- **lane**（各行）来自代码中 `group` 与 `lane` 的名字，例如 `producer / tma`、`consumer / wgmma`。每个 `range` 是一段条形，鼠标悬停可以看到它的名字与周期区间。
+- **横轴**是 SM 的原始周期数（`clock64()`），每个 CTA 从零开始。
+- **箭头**是代码中声明的 `dag` 边（例如 producer 的 `arrive` → consumer 的 `wait`），每出现一次画一条。从箭头可以读出交接的延迟，以及 consumer 是否在空等。
+- 缩放与平移只在水平方向进行。
+
+需要留意的现象：
+
+- `wgmma` 这一行上有空隙：consumer 停下来等待 TMA 加载。
+- 相邻迭代之间的 `dag` 箭头没有重叠：没有形成流水。
+- 某一行明显比其他行长：负载不均衡。
diff --git a/docs/user-guide/development.md b/docs/user-guide/development.md
new file mode 100644
index 00000000..40c585e5
--- /dev/null
+++ b/docs/user-guide/development.md
@@ -0,0 +1,190 @@
+# Development Guide { #development-guide }
+
+## Getting the code { #get-code }
+
+```bash
+git clone https://github.com/<your-name>/TileOPs
+cd TileOPs
+git remote add upstream https://github.com/tile-ai/TileOPs
+git fetch upstream
+git switch -c <branch> upstream/main
+```
+
+`<your-name>` is the GitHub account that holds the fork of
+[tile-ai/TileOPs](https://github.com/tile-ai/TileOPs). A development branch starts
+from the latest upstream `main`, and the upstream repository is recorded as
+`upstream`.
+
+## Setting up an environment { #setup }
+
+### The dev image { #docker }
+
+```bash
+docker run --rm -it --gpus all \
+  -v "$(pwd)":/workspace -w /workspace \
+  ghcr.io/tile-ai/tileops-runner:<tag>
+
+# inside the container
+pip install -e . --no-deps --no-build-isolation
+```
+
+- The dev images are published at
+  [ghcr.io/tile-ai/tileops-runner](https://github.com/tile-ai/TileOPs/pkgs/container/tileops-runner);
+  development uses a tag ending in `-dev`.
+- The image carries CUDA, PyTorch, TileLang and the test tools, and comes from the
+  same build as the CI's GPU runners.
+- `--no-deps` skips resolving and installing dependencies, and uses the ones the
+  image already has.
+
+The dev image does not include pre-commit. Run the following on the host, or in
+any other environment with Python:
+
+```bash
+pip install pre-commit
+pre-commit install
+```
+
+### A local environment { #local }
+
+```bash
+pip install -e '.[dev]' -c constraints.txt
+pre-commit install
+```
+
+- The supported combination of Python, PyTorch, CUDA, GPU architecture and
+  TileLang is the one listed under Prerequisites in the TileOPs
+  [README](https://github.com/tile-ai/TileOPs#installation).
+- `-c constraints.txt` applies the repository's dependency constraints, so the
+  local dependencies match the combination CI validates.
+
+### Checking the environment { #verify }
+
+```bash
+python -m pytest -q tests -m smoke
+```
+
+Each kernel under test is compiled and cached on its first call, and later
+identical calls reuse the cache; see [Kernel compilation and caching](#compile).
+
+## Making a change { #change }
+
+### The relevant spec and design docs { #design-first }
+
+The design docs and the manifest define the constraints on ops, kernels and
+tests. When a change to the implementation affects those constraints, the same PR
+updates the spec.
+
+| Change | Documents |
+| --- | --- |
+| Adding an op | [Adding a new op](../new-op.md) |
+| An existing op or kernel | [Op Interfaces](../design/ops-design.md), [Slot Rules](../design/op-slot-rules.md) |
+| A spec in the manifest | [Spec Fields](manifest/writing.md), [Manifest](../design/manifest.md) |
+| Tests | [Testing](../design/testing.md), [Layer Boundaries](../design/layer-boundaries.md) |
+| An op's docstring | [Docstrings](https://github.com/tile-ai/TileOPs/blob/main/docs/development.md#docstrings) |
+
+### Kernel compilation and caching { #compile }
+
+```python
+op = GemmFwdOp()
+d = op(a, b)   # first call: compiles the kernel and caches it
+d = op(a, b)   # the same call: reuses the cache
+```
+
+- Installing TileOPs compiles no kernels. TileLang compiles a kernel the first time
+  an op is called.
+- An op looks up the entry serving a call by the whole call; what is compiled is
+  cached by build identity, and calls that select the same implementation class
+  with the same build identity share one entry. TileLang compiles when no entry
+  matches. Build identity is defined in
+  [How an op selects a kernel](dispatch/index.md).
+- Under an editable install, changes to Python source take effect directly; a
+  change to dependencies or build configuration needs the install command run
+  again.
+
+## Running the tests { #tests }
+
+| Command | Tests included |
+| --- | --- |
+| `python -m pytest -q tests -m smoke` | `smoke`: the critical path |
+| `python -m pytest -q tests -m "smoke or full"` | adds `full`: standard correctness coverage |
+| `python -m pytest -q tests -m "smoke or full or nightly"` | adds `nightly`: exhaustive and long-running cases |
+| `python -m pytest -q tests` | every test, with no marker filter |
+
+A single test file runs with `python -m pytest -q <test-file>`.
+
+These checks need no GPU:
+
+```bash
+python -m pytest -q tests/test_validate_manifest.py   # manifest spec validation
+python -m pytest -q benchmarks/tests                  # benchmark infrastructure tests
+pre-commit run --all-files                            # lint, the same as CI's pre-commit check
+```
+
+## Running benchmarks { #bench }
+
+```bash
+PIP_NO_BUILD_ISOLATION=1 pip install -e '.[dev,bench]' -c constraints.txt
+python -m pytest -q <bench-file>
+```
+
+- A PR that changes a kernel or an op includes benchmark results, compared against
+  an implementation outside TileOPs.
+- The baseline libraries install through the `bench` extra. The dev image installs
+  them at build time too; a baseline that fails to install there does not stop the
+  build, and `sgl-kernel` is not in the image. A baseline missing from the
+  container is installed at the version the repository declares.
+- How the numbers are timed: [How a benchmark is timed](../timing.md).
+
+## Opening a PR { #pr }
+
+### Title { #pr-title }
+
+| Format | Example |
+| --- | --- |
+| `[Type] <description>` | `[Doc] Fix the install command in the README` |
+| `[Type][Scope] <description>` | `[BugFix][Elementwise] Build the floored tiers at every tuned fold width` |
+| `[Type][foundry][Scope] <description>` | `[Perf][foundry][Elementwise] Take the floored tier's reciprocal from one MUFU instruction` |
+
+- CI checks the PR title's format. The values of `Type` are defined in
+  [`.claude/conventions/types.sh`](https://github.com/tile-ai/TileOPs/blob/main/.claude/conventions/types.sh).
+- `foundry` marks a PR whose kernels were generated by
+  [TileFoundry](https://github.com/tile-ai/TileFoundry).
+
+### Description { #pr-body }
+
+The PR description follows the
+[PR template](https://github.com/tile-ai/TileOPs/blob/main/.github/PULL_REQUEST_TEMPLATE.md).
+When a change touches `tests/`, the description includes the change in the test
+count:
+
+```bash
+python scripts/test_node_delta.py --base upstream/main
+```
+
+### CI { #ci }
+
+- A draft PR runs the PR title check and the manifest statistics only, and skips
+  the CPU checks and GPU smoke tests below. They run once the PR is marked
+  **Ready for review**.
+- The CPU checks are pre-commit, gitleaks, manifest validation, actionlint, the
+  compile contract check and the packaging check, plus the benchmark contract tests
+  when a change touches the benchmarks.
+- The GPU smoke tests run once pre-commit, gitleaks and actionlint pass, scoped by
+  the files the PR changes. A PR that changes no Python, manifest or native source
+  skips them.
+
+## FAQ { #faq }
+
+### A local install fails while re-resolving CUDA or TileLang { #faq-rebuild }
+
+```bash
+PIP_NO_BUILD_ISOLATION=1 pip install -e '.[dev]' -c constraints.txt
+```
+
+With CUDA and TileLang already installed, turning off build isolation lets the
+build use the installed versions.
+
+### Checks that run without an SM90 GPU { #faq-no-gpu }
+
+The three GPU-free checks under [Running the tests](#tests). The tests that need a
+GPU run in the PR's CI.
diff --git a/docs/user-guide/development.zh.md b/docs/user-guide/development.zh.md
new file mode 100644
index 00000000..3cff28ad
--- /dev/null
+++ b/docs/user-guide/development.zh.md
@@ -0,0 +1,152 @@
+# 开发指南 { #development-guide }
+
+## 获取代码 { #get-code }
+
+```bash
+git clone https://github.com/<your-name>/TileOPs
+cd TileOPs
+git remote add upstream https://github.com/tile-ai/TileOPs
+git fetch upstream
+git switch -c <branch> upstream/main
+```
+
+`<your-name>` 是 fork [tile-ai/TileOPs](https://github.com/tile-ai/TileOPs) 所在的 GitHub 账号。开发分支从上游最新的 `main` 创建，上游仓库记录为 `upstream`。
+
+## 搭建开发环境 { #setup }
+
+### 使用 dev 镜像 { #docker }
+
+```bash
+docker run --rm -it --gpus all \
+  -v "$(pwd)":/workspace -w /workspace \
+  ghcr.io/tile-ai/tileops-runner:<tag>
+
+# 以下命令在容器内执行
+pip install -e . --no-deps --no-build-isolation
+```
+
+- dev 镜像发布在 [ghcr.io/tile-ai/tileops-runner](https://github.com/tile-ai/TileOPs/pkgs/container/tileops-runner)，开发使用以 `-dev` 结尾的 tag。
+- 镜像包含 CUDA、PyTorch、TileLang 与测试工具，与 CI 的 GPU runner 出自同一构建流程。
+- `--no-deps` 跳过依赖的解析与安装，直接使用镜像中已有的依赖。
+
+dev 镜像不包含 pre-commit。以下命令在宿主机或另一个已安装 Python 的环境中执行：
+
+```bash
+pip install pre-commit
+pre-commit install
+```
+
+### 使用本地环境 { #local }
+
+```bash
+pip install -e '.[dev]' -c constraints.txt
+pre-commit install
+```
+
+- 支持的 Python、PyTorch、CUDA、GPU 架构与 TileLang 版本组合，以 TileOPs 仓库 [README](https://github.com/tile-ai/TileOPs#installation) 的 Prerequisites 为准。
+- `-c constraints.txt` 使用仓库提供的依赖约束，使本地依赖与 CI 验证的组合一致。
+
+### 确认环境可用 { #verify }
+
+```bash
+python -m pytest -q tests -m smoke
+```
+
+被测的 kernel 在首次调用时编译并写入缓存，之后相同的调用复用缓存，见 [kernel 的编译与缓存](#compile)。
+
+## 修改代码 { #change }
+
+### 相关的 spec 与设计文档 { #design-first }
+
+设计文档与 manifest 定义 op、kernel 与测试的约束。实现的变更影响这些约束时，同一个 PR 同步更新对应的 spec。
+
+| 改动内容 | 相关文档 |
+| --- | --- |
+| 新增 op | [添加新 op](../new-op.md) |
+| 修改已有的 op 或 kernel | [Op 接口](../design/ops-design.md)、[Slot 规则](../design/op-slot-rules.md) |
+| 修改 manifest 中的 spec | [写一个 spec](manifest/writing.md)、[manifest 规范](../design/manifest.md) |
+| 修改测试 | [测试](../design/testing.md)、[层间边界](../design/layer-boundaries.md) |
+| 修改 op 的 docstring | [docstring 的写法](https://github.com/tile-ai/TileOPs/blob/main/docs/development.md#docstrings) |
+
+### kernel 的编译与缓存 { #compile }
+
+```python
+op = GemmFwdOp()
+d = op(a, b)   # 首次调用：编译 kernel 并写入缓存
+d = op(a, b)   # 相同的调用：复用缓存
+```
+
+- 安装 TileOPs 时不编译 kernel。kernel 在 op 首次被调用时由 TileLang 编译。
+- op 按完整的调用信息查找服务这次调用的 entry；编译结果按 build identity 缓存，选中同一实现类且 build identity 相同的调用共用同一个 entry。没有对应的 entry 时，TileLang 执行编译。build identity 的定义见 [op 如何选择 kernel](dispatch/index.md)。
+- editable 安装下，Python 源码的改动直接生效；依赖或构建配置变更后，需要重新执行安装命令。
+
+## 运行测试 { #tests }
+
+| 命令 | 包含的测试 |
+| --- | --- |
+| `python -m pytest -q tests -m smoke` | `smoke`：关键路径 |
+| `python -m pytest -q tests -m "smoke or full"` | 另含 `full`：标准的正确性覆盖 |
+| `python -m pytest -q tests -m "smoke or full or nightly"` | 另含 `nightly`：穷举与耗时长的用例 |
+| `python -m pytest -q tests` | 全部测试，不按 marker 过滤 |
+
+单个测试文件以 `python -m pytest -q <test-file>` 运行。
+
+以下检查不需要 GPU：
+
+```bash
+python -m pytest -q tests/test_validate_manifest.py   # manifest spec 校验
+python -m pytest -q benchmarks/tests                  # benchmark 基础设施测试
+pre-commit run --all-files                            # 代码检查，与 CI 中的 pre-commit 一项相同
+```
+
+## 运行 benchmark { #bench }
+
+```bash
+PIP_NO_BUILD_ISOLATION=1 pip install -e '.[dev,bench]' -c constraints.txt
+python -m pytest -q <bench-file>
+```
+
+- 改动 kernel 或 op 的 PR 需附上 benchmark 结果，对比对象是 TileOPs 以外的实现。
+- benchmark 的对比库通过 `bench` extra 安装。dev 镜像在构建时也会安装这些库，其中个别库安装失败不会中止构建，`sgl-kernel` 则不在镜像中；容器内缺少的库按仓库声明的版本另行安装。
+- 计时方式见 [benchmark 的计时方法](../timing.md)。
+
+## 提交 PR { #pr }
+
+### 标题 { #pr-title }
+
+| 格式 | 例子 |
+| --- | --- |
+| `[Type] <description>` | `[Doc] Fix the install command in the README` |
+| `[Type][Scope] <description>` | `[BugFix][Elementwise] Build the floored tiers at every tuned fold width` |
+| `[Type][foundry][Scope] <description>` | `[Perf][foundry][Elementwise] Take the floored tier's reciprocal from one MUFU instruction` |
+
+- CI 检查 PR 标题的格式。`Type` 的取值定义在 [`.claude/conventions/types.sh`](https://github.com/tile-ai/TileOPs/blob/main/.claude/conventions/types.sh) 中。
+- `foundry` 标记 kernel 由 [TileFoundry](https://github.com/tile-ai/TileFoundry) 生成的 PR。
+
+### 描述 { #pr-body }
+
+PR 描述按 [PR 模板](https://github.com/tile-ai/TileOPs/blob/main/.github/PULL_REQUEST_TEMPLATE.md)填写。改动涉及 `tests/` 时，描述中附上测试用例数量的变化：
+
+```bash
+python scripts/test_node_delta.py --base upstream/main
+```
+
+### CI { #ci }
+
+- Draft PR 只运行 PR 标题检查与 manifest 统计，跳过下面的 CPU 检查与 GPU smoke 测试。PR 转为 **Ready for review** 后，这些检查才会运行。
+- CPU 检查包括 pre-commit、gitleaks、manifest 校验、actionlint、编译约定检查与打包检查；改动涉及 benchmark 时，另含 benchmark 约定测试。
+- GPU smoke 测试在 pre-commit、gitleaks 与 actionlint 通过后运行，测试范围由改动的文件决定。没有改动 Python、manifest 或原生代码的 PR 跳过 GPU smoke。
+
+## 常见问题 { #faq }
+
+### 本地安装在重新解析 CUDA 或 TileLang 时失败 { #faq-rebuild }
+
+```bash
+PIP_NO_BUILD_ISOLATION=1 pip install -e '.[dev]' -c constraints.txt
+```
+
+本机已安装 CUDA 与 TileLang 时，关闭构建隔离，构建过程直接使用已安装的版本。
+
+### 没有 SM90 GPU 时可运行的检查 { #faq-no-gpu }
+
+[运行测试](#tests)一节中不需要 GPU 的三项检查。需要 GPU 的测试由 PR 上的 CI 运行。
diff --git a/docs/user-guide/index.md b/docs/user-guide/index.md
index 8c82f1d5..bb973373 100644
--- a/docs/user-guide/index.md
+++ b/docs/user-guide/index.md
@@ -1,10 +1,30 @@
 # User Guide
 
-| Page | Contents |
-| --- | --- |
-| [Reading and writing the manifest](manifest/index.md) | how the system uses a spec, the concepts used to describe a spec, and how to write one |
-| [How an op selects a kernel](dispatch/index.md) | how an op selects a kernel, how to add a kernel, and how a backend joins |
-| [Adding a new op](../new-op.md) | the six steps from a spec to `status: implemented` |
-| [Bringing an op into torch.compile](../torch-compile.md) | what an op looks like inside a compiled graph, and the conventions a caller follows |
-| [How a benchmark is timed](../timing.md) | how the numbers on the Benchmarks pages are measured |
-| [Adding a hardware backend](../backends.md) | serving the ops on one class of devices with your own kernels |
+## Getting started
+
+- [Development guide](development.md)
+  Getting the code, setting up the dev image or a local environment, running the
+  tests, and opening a PR.
+
+## Specs and ops
+
+- [Reading and writing the manifest](manifest/index.md)
+  How the system uses a spec, the concepts a spec is described in, and how to
+  write one.
+- [Adding a new op](../new-op.md)
+  The six steps from a spec to `status: implemented`.
+
+## Kernels and hardware
+
+- [How an op selects a kernel](dispatch/index.md)
+  How an op selects a kernel, how to add a kernel, and how a backend joins.
+- [Adding a hardware backend](../backends.md)
+  Serving the ops on one class of devices with your own kernels.
+
+## Integration and performance
+
+- [Bringing an op into torch.compile](../torch-compile.md)
+  What an op looks like inside a compiled graph, and the conventions a caller
+  follows.
+- [How a benchmark is timed](../timing.md)
+  How the numbers on the Benchmarks pages are measured.
diff --git a/docs/user-guide/index.zh.md b/docs/user-guide/index.zh.md
index ac2a3fd6..3e18c353 100644
--- a/docs/user-guide/index.zh.md
+++ b/docs/user-guide/index.zh.md
@@ -1,10 +1,27 @@
-# 使用指南
-
-| 文档 | 内容 |
-| --- | --- |
-| [读写 manifest](manifest/index.md) | 系统如何使用 spec、描述 spec 所用的概念，以及如何写一份 spec |
-| [op 如何选择 kernel](dispatch/index.md) | op 如何选中 kernel、如何新增 kernel，以及 backend 如何接入 |
-| [添加新 op](../new-op.md) | 从一份 spec 到 `status: implemented` 的六步 |
-| [接入 torch.compile](../torch-compile.md) | op 在编译图中的形态，以及调用时的约定 |
-| [benchmark 的计时方法](../timing.md) | Benchmarks 页上的数字如何测得 |
-| [接入新硬件 backend](../backends.md) | 用自己的 kernel 接管某一类设备上的 op |
+# 用户指南
+
+## 开始开发
+
+- [开发指南](development.md)
+  获取代码，搭建 dev 镜像或本地环境，运行测试，提交 PR。
+
+## spec 与 op
+
+- [读写 manifest](manifest/index.md)
+  系统如何使用 spec，描述 spec 所用的概念，以及 spec 的写法。
+- [添加新 op](../new-op.md)
+  从一份 spec 到 `status: implemented` 的六个步骤。
+
+## kernel 与硬件
+
+- [op 如何选择 kernel](dispatch/index.md)
+  op 选中 kernel 的过程，新增 kernel 的方法，以及 backend 的接入方式。
+- [接入新硬件 backend](../backends.md)
+  在某一类设备上，用自己的 kernel 实现 op。
+
+## 集成与性能
+
+- [接入 torch.compile](../torch-compile.md)
+  op 在编译图中的形态，以及调用时的约定。
+- [benchmark 的计时方法](../timing.md)
+  性能数据页上的数字如何测得。
diff --git a/hooks.py b/hooks.py
index 0c144ff2..9b7b7ded 100644
--- a/hooks.py
+++ b/hooks.py
@@ -10,6 +10,12 @@
 * mkdocs-static-i18n serves the default-language page where a locale has no
   translation; `on_page_markdown` prepends a notice, so a fallback page reads as
   a translation still to come rather than a broken one.
+* A blog post's first paragraph is its subtitle; `on_page_content` marks it for
+  extra.css, so the post's Markdown carries no styling.
+* A section index lists its pages as a plain Markdown list; `on_page_content`
+  marks those lists for extra.css, which draws each item as a card.
+* A page merged into another leaves its old URL behind; `on_post_build` writes
+  a redirect there, so published links keep working.
 """
 from __future__ import annotations
 
@@ -69,9 +75,9 @@ def on_page_markdown(markdown, page, config, files):
 # not produce is left out; one it produced that is not listed here is appended.
 _BENCH_ORDER = [
     "index.md", "reading.md",
-    "elementwise.md", "reduction.md", "normalization.md", "quantization.md",
-    "gemm.md", "conv-pool.md", "moe.md", "sampling.md", "rope.md",
-    "attention.md", "linear-attention.md", "ssm.md", "other.md",
+    "elementwise.md", "rope.md", "reduction.md", "normalization.md",
+    "conv-pool.md", "gemm.md", "quantization.md", "attention.md", "moe.md",
+    "sampling.md", "linear-attention.md", "ssm.md", "other.md",
 ]
 
 
@@ -89,3 +95,63 @@ def on_config(config):
         if isinstance(section, dict) and "Benchmarks" in section:
             section["Benchmarks"] = entries
     return config
+
+
+_FIRST_PARAGRAPH = re.compile(r"(</h1>\s*)<p>")
+
+# Section indexes whose page lists render as cards. Each item is a link followed
+# by a one-line description on the next line.
+_CARD_INDEXES = {
+    "user-guide/index.md", "user-guide/index.zh.md", "design/index.md",
+}
+
+
+def on_page_content(html, page, config, files):
+    """Mark a blog post's subtitle, and the page lists on a section index."""
+    src = page.file.src_path.replace("\\", "/")
+    name = src.split("/")[-1]
+    if src in _CARD_INDEXES:
+        return html.replace("<ul>", '<ul class="page-cards">')
+    if not src.startswith("blog/") or name.startswith("index."):
+        return html
+    return _FIRST_PARAGRAPH.sub(r'\1<p class="post-subtitle">', html, count=1)
+
+
+# Old page URL -> where its content lives now, relative to the old URL. The
+# redirect is written into each locale's build.
+_REDIRECTS = {
+    "api/topk/": "../sampling/#top-k-selection",
+}
+
+_REDIRECT_PAGE = (
+    '<!doctype html><html><head><meta charset="utf-8">'
+    '<meta http-equiv="refresh" content="0; url={to}">'
+    '<link rel="canonical" href="{to}"><title>Redirecting</title></head>'
+    '<body><a href="{to}">{to}</a></body></html>\n'
+)
+
+
+def _locale_dirs(config):
+    """The site root, then one subdirectory per other locale i18n builds.
+
+    This hook runs before i18n builds the other locales, so their directories
+    are named from the plugin's config rather than found on disk.
+    """
+    dirs = [""]
+    i18n = config["plugins"].get("i18n")
+    for lang in getattr(i18n, "config", {}).get("languages", []):
+        if lang.build and not lang.default:
+            dirs.append(lang.locale)
+    return dirs
+
+
+def on_post_build(config):
+    """Write a redirect page at each old URL `_REDIRECTS` lists, per locale."""
+    for locale in _locale_dirs(config):
+        for old, to in _REDIRECTS.items():
+            path = os.path.join(config["site_dir"], locale, old, "index.html")
+            if os.path.exists(path):
+                continue
+            os.makedirs(os.path.dirname(path), exist_ok=True)
+            with open(path, "w", encoding="utf-8") as f:
+                f.write(_REDIRECT_PAGE.format(to=to))
diff --git a/mkdocs.yml b/mkdocs.yml
index 97ac87ba..73f60e20 100644
--- a/mkdocs.yml
+++ b/mkdocs.yml
@@ -29,8 +29,13 @@ theme:
 
 nav:
   - Home: index.md
+  - Blog:
+    - blog/index.md
   - User Guide:
+    # In the order the index groups them: getting started, specs and ops,
+    # kernels and hardware, integration and performance.
     - user-guide/index.md
+    - Development Guide: user-guide/development.md
     - Reading and Writing Specs:
       - user-guide/manifest/index.md
       - Concepts: user-guide/manifest/concepts.md
@@ -38,39 +43,40 @@ nav:
       - Extensions: user-guide/manifest/extensions.md
       - Calls and Validation: user-guide/manifest/calls.md
       - Examples: user-guide/manifest/examples.md
+    - Adding an Op: new-op.md
     - Kernel Dispatch:
       - user-guide/dispatch/index.md
       - Adding a Kernel: user-guide/dispatch/writing.md
       - Backend Integration: user-guide/dispatch/backends.md
-    - Adding an Op: new-op.md
-    - Timing Benchmarks: timing.md
-    - Using torch.compile: torch-compile.md
     - Adding a Backend: backends.md
-  # In `tileops._FAMILIES` order — simple to composite: the pointwise transforms,
-  # then the axis reductions and the normalizations built on them, then the matmul
-  # and the expert routing over it, then the windowed and spectral transforms, then
-  # the sequence-model kernels built on all of the above. Pages are not one per
-  # family. The FP8 lightning indexer is a section of the Attention page, where its
-  # single op is used, and the sequence-modeling family is two pages, mHC and Engram,
-  # being two unrelated algorithms. Top-k keeps its page before GEMM though its op is
-  # exported from `tileops.attention`. Trace comes last — a tool rather than an op.
+    - Using torch.compile: torch-compile.md
+    - Timing Benchmarks: timing.md
+  # In the Benchmarks pages' order: the pointwise transforms with the positional
+  # rotation beside them, the axis reductions and the normalizations built on
+  # them, the windowed kernels, the matmul and the quantization around it, then
+  # attention, the expert routing after it and sampling at the end of a decode
+  # step, then the sequence-mixing kernels. The pages Benchmarks folds into Other
+  # follow: FFT, mHC and Engram. Top-k shares the Sampling page, though its op is
+  # exported from `tileops.attention` and its benchmark sits on Attention. Pages are not one per family. The FP8
+  # lightning indexer is a section of the Attention page, where its single op is
+  # used, and the sequence-modeling family is two pages, mHC and Engram, being two
+  # unrelated algorithms. Trace comes last — a tool rather than an op.
   - API Reference:
     - api/index.md
     - Elementwise: api/elementwise.md
+    - RoPE: api/rope.md
     - Reduction: api/reduction.md
     - Normalization: api/normalization.md
-    - Quantization: api/quantization.md
-    - Top-k: api/topk.md
-    - GEMM: api/linear-algebra.md
     - Pooling: api/pool.md
     - Convolution: api/convolution.md
-    - FFT: api/fft.md
-    - MoE: api/moe.md
-    - Sampling: api/sampling.md
-    - RoPE: api/rope.md
+    - GEMM: api/linear-algebra.md
+    - Quantization & Dequantization: api/quantization.md
     - Attention: api/attention.md
+    - MoE: api/moe.md
+    - Top-k & Sampling: api/sampling.md
     - Linear Attention: api/linear-attention.md
     - Mamba: api/mamba.md
+    - FFT: api/fft.md
     - mHC: api/mhc.md
     - Engram: api/engram.md
     - Trace: api/trace.md
@@ -83,16 +89,18 @@ nav:
       - performance-guides/memory-bound/index.md
       - Global Memory Access: performance-guides/memory-bound/global-memory-access.md
       - Shared Memory Access: performance-guides/memory-bound/shared-memory-access.md
+  # In the order the index groups them: the system's structure, the spec, building
+  # an op against it, then how an op is accepted.
   - Design:
     - design/index.md
     - Architecture: design/architecture.md
+    - Layer Boundaries: design/layer-boundaries.md
     - Manifest: design/manifest.md
     - Op Interfaces: design/ops-design.md
     - Op Interface Reference: design/ops-design-reference.md
     - Slot Rules: design/op-slot-rules.md
-    - Roofline: design/roofline.md
     - Testing: design/testing.md
-    - Layer Boundaries: design/layer-boundaries.md
+    - Roofline: design/roofline.md
 
 markdown_extensions:
   # Docstrings and design docs carry formulas. `generic: true` emits the plain
@@ -160,8 +168,10 @@ plugins:
           build: true
           nav_translations:
             Home: 首页
+            Blog: 博客
             Design: 设计文档
-            User Guide: 使用指南
+            User Guide: 用户指南
+            Development Guide: 开发指南
             Reading and Writing Specs: 读写 manifest
             Concepts: 概念
             Spec Fields: 写一个 spec
@@ -178,6 +188,7 @@ plugins:
             Architecture: 架构
             Manifest: manifest 规范
             Op Interfaces: op 接口
+            Op Interface Reference: op 接口参考
             Slot Rules: Slot 规则
             Roofline: Roofline
             Testing: 测试
@@ -188,6 +199,7 @@ plugins:
             Normalization: Normalization
             GEMM: GEMM
             Quantization: Quantization
+            Quantization & Dequantization: Quantization & Dequantization
             Attention: Attention
             RoPE: RoPE
             Convolution: Convolution
@@ -199,7 +211,7 @@ plugins:
             MoE: MoE
             Sampling: Sampling
             Engram: Engram
-            Top-k: Top-k
+            Top-k & Sampling: Top-k & Sampling
             Trace: Trace
             Conv & Pool: Conv & Pool
             SSM: SSM
diff --git a/scripts/gen_bench_pages.py b/scripts/gen_bench_pages.py
index 8f85659d..67018301 100644
--- a/scripts/gen_bench_pages.py
+++ b/scripts/gen_bench_pages.py
@@ -2,8 +2,8 @@
 """Render the Benchmarks section from a nightly benchmark XML snapshot.
 
 Output is one overview page, one page explaining the numbers, and the data
-pages of `DATA_PAGES`, in the order the API Reference nav lists the same
-families. `hooks.py` puts them into the site nav in that order.
+pages of `DATA_PAGES`, in the order that list gives. `hooks.py` puts them into
+the site nav in that order.
 
 These pages answer one question per workload: **how does TileOPs compare to the
 fastest other implementation of the same op on that workload?**
@@ -96,7 +96,7 @@ def tier_of(tag: str) -> str:
     "ssm": "SSM", "scan": "Scan", "normalization": "Normalization",
     "moe": "MoE", "linear_algebra": "GEMM", "reduction": "Reduction",
     "elementwise": "Elementwise", "convolution": "Convolution", "pool": "Pooling",
-    "quantization": "Quantization", "sampling": "Sampling", "positional": "RoPE",
+    "quantization": "Quantization & Dequantization", "sampling": "Sampling", "positional": "RoPE",
     "fft": "FFT", "mhc": "mHC", "engram": "Engram", "topk": "Top-k",
     "other": "Other",
 }
@@ -142,24 +142,24 @@ def api_op_order(api_dir: str = API_DIR,
         print(f"warning: no op order read from {api_dir}; every Benchmarks page "
               f"ranks its ops by verdict instead", file=sys.stderr)
     return order
-# (slug, page title, families in display order), in the order the API Reference
-# nav lists the same families — pointwise, then the reductions and the
-# normalizations built on them, then quantization, then the matrix multiply and
-# the expert routing over it, then the sampling, then the positional rotation,
-# then the sequence-mixing kernels built on all of the above. A page is one
-# family except where too few ops carry one: `Conv & Pool` is two, `Other` the
-# rest.
+# (slug, page title, families in display order). Pointwise first, with the
+# positional rotation beside it, then the reductions and the normalizations
+# built on them, the windowed kernels, the matrix multiply and the quantization
+# around it, then attention and the expert routing that follows it in a
+# transformer layer, sampling at the end of a decode step, then the
+# sequence-mixing kernels. A page is one family except where too few ops carry
+# one: `Conv & Pool` is two, `Other` the rest.
 DATA_PAGES = [
     ("elementwise", "Elementwise", ["elementwise"]),
+    ("rope", "RoPE", ["positional"]),
     ("reduction", "Reduction", ["reduction"]),
     ("normalization", "Normalization", ["normalization"]),
-    ("quantization", "Quantization", ["quantization"]),
-    ("gemm", "GEMM", ["linear_algebra"]),
     ("conv-pool", "Conv & Pool", ["pool", "convolution"]),
+    ("gemm", "GEMM", ["linear_algebra"]),
+    ("quantization", "Quantization & Dequantization", ["quantization"]),
+    ("attention", "Attention", ["attention"]),
     ("moe", "MoE", ["moe"]),
     ("sampling", "Sampling", ["sampling"]),
-    ("rope", "RoPE", ["positional"]),
-    ("attention", "Attention", ["attention"]),
     ("linear-attention", "Linear Attention", ["linear_attention"]),
     ("ssm", "SSM", ["ssm"]),
     ("other", "Other", ["topk", "fft", "mhc", "engram", "scan", "other"]),