diff --git a/_extra/viz-base.js b/_extra/viz-base.js index f3339bc2..ce3e1a3f 100644 --- a/_extra/viz-base.js +++ b/_extra/viz-base.js @@ -3,6 +3,17 @@ document.addEventListener('DOMContentLoaded', function() { var p = new URLSearchParams(location.search); if (p.has('notitle')) document.body.classList.add('notitle'); + // Directly opened demo pages are copied outside Sphinx templates, so they do + // not receive html_js_files. Load the same language switch used by book pages. + if (window.parent === window && !p.has('notitle')) { + var langSwitchPath = window.location.pathname.indexOf('/zh/demo/') >= 0 + ? '../../_static/lang-switch.js' + : '../_static/lang-switch.js'; + var s = document.createElement('script'); + s.src = new URL(langSwitchPath, window.location.href).href; + document.head.appendChild(s); + } + // Forward arrow keys to parent (reveal.js) when embedded if (window.parent !== window) { document.addEventListener('keydown', function(e) { diff --git a/_extra_zh/zh/_static/tirx-layout-demo/index.html b/_extra_zh/zh/_static/tirx-layout-demo/index.html new file mode 100644 index 00000000..ae2ec85e --- /dev/null +++ b/_extra_zh/zh/_static/tirx-layout-demo/index.html @@ -0,0 +1,167 @@ + + + +
+ + +TileLayout (shard / replica / offset) maps logical tensor elements to physical threads…or
…+ * + * Token classes: + * .kw — keyword (if, def, for, while, with, return, True, False) + * .fn — callable (function/method calls — the prominent highlight) + * .str — string literal + * .num — number literal + * .cmt — comment + * .op — operator (+, -, ==, =, :) + * .typ — type name + * .dec — decorator (@) + * + * Convention: only wrap the callable name in .fn, NOT the namespace prefix. + * ✓ Tx.copy_async(...) + * ✗ Tx.copy_async(...) + */ + +/* ── Dark theme (dark bg, bright tokens) ────────────── */ +.code-dark { + background: #0f172a; color: #e2e8f0; + font-family: 'SF Mono','Fira Code',monospace; + font-size: 12px; line-height: 1.6; + border-radius: 8px; padding: 12px 16px; + overflow-x: auto; +} +.code-dark .kw { color: #c084fc; } /* purple — keywords */ +.code-dark .fn { color: #60a5fa; } /* blue — callables (prominent) */ +.code-dark .str { color: #34d399; } /* green — strings */ +.code-dark .num { color: #fbbf24; } /* yellow — numbers */ +.code-dark .cmt { color: #64748b; font-style: italic; } /* gray — comments */ +.code-dark .op { color: #94a3b8; } /* slate — operators */ +.code-dark .typ { color: #7dd3fc; } /* cyan — types */ +.code-dark .dec { color: #c084fc; } /* purple — decorators */ + +/* ── Light theme (white bg, muted tokens) ───────────── */ +.code-light { + background: #f8fafc; color: #1e293b; + font-family: 'SF Mono','Fira Code',monospace; + font-size: 12px; line-height: 1.6; + border-radius: 8px; padding: 12px 16px; + border: 1px solid #e2e8f0; + overflow-x: auto; +} +.code-light .kw { color: #7c3aed; } /* purple — keywords */ +.code-light .fn { color: #2563eb; } /* blue — callables (prominent) */ +.code-light .str { color: #059669; } /* green — strings */ +.code-light .num { color: #d97706; } /* amber — numbers */ +.code-light .cmt { color: #6b7280; font-style: italic; } /* gray — comments */ +.code-light .op { color: #64748b; } /* slate — operators */ +.code-light .typ { color: #0369a1; } /* cyan — types */ +.code-light .dec { color: #7c3aed; } /* purple — decorators */ + +/* ── Code block panel (used by createCodeBlock) ────── + * + * Structure: + *
— code container (scrollable) + * (regions and lines rendered by createCodeBlock) + *+ *
works too.
+ */
+
+.cb-panel {
+ border-radius: 8px;
+ overflow: hidden;
+}
+.cb-panel > .code-dark,
+.cb-panel > .code-light {
+ border-radius: 0;
+}
+
+/* Header bar */
+.cb-panel .cb-header {
+ padding: 6px 14px;
+ font-size: 11px;
+ font-weight: 600;
+ font-family: -apple-system, BlinkMacSystemFont, 'Segoe UI', sans-serif;
+}
+.cb-panel.dark .cb-header {
+ background: #1e293b;
+ color: #94a3b8;
+ border-bottom: 1px solid rgba(255,255,255,0.06);
+}
+.cb-panel.light .cb-header {
+ background: #f1f5f9;
+ color: #64748b;
+ border-bottom: 1px solid #e2e8f0;
+}
+
+/* Dark scrollbar */
+.code-dark {
+ scrollbar-width: thin;
+ scrollbar-color: rgba(255,255,255,0.15) transparent;
+}
+.code-dark::-webkit-scrollbar { width: 6px; height: 6px; }
+.code-dark::-webkit-scrollbar-track { background: transparent; }
+.code-dark::-webkit-scrollbar-thumb { background: rgba(255,255,255,0.15); border-radius: 3px; }
+.code-dark::-webkit-scrollbar-thumb:hover { background: rgba(255,255,255,0.25); }
+
+/* Light scrollbar */
+.code-light {
+ scrollbar-width: thin;
+ scrollbar-color: rgba(0,0,0,0.15) transparent;
+}
+.code-light::-webkit-scrollbar { width: 6px; height: 6px; }
+.code-light::-webkit-scrollbar-track { background: transparent; }
+.code-light::-webkit-scrollbar-thumb { background: rgba(0,0,0,0.12); border-radius: 3px; }
+.code-light::-webkit-scrollbar-thumb:hover { background: rgba(0,0,0,0.2); }
+
+/* Regions */
+.cb-region {
+ border-left: 3px solid transparent;
+ padding: 0.35rem 0 0.35rem 0.85rem;
+ margin: 0.15rem 0;
+ border-radius: 2px;
+ cursor: pointer;
+ transition: background 0.2s;
+}
+.code-dark .cb-region:hover { background: rgba(255,255,255,0.05); }
+.code-dark .cb-region.active { background: rgba(255,255,255,0.07); }
+.code-light .cb-region:hover { background: rgba(0,0,0,0.04); }
+.code-light .cb-region.active { background: rgba(0,0,0,0.06); }
+
+/* Lines */
+.cb-line {
+ display: block;
+ white-space: pre;
+}
+.cb-line.cb-focus {
+ margin-left: -0.85rem;
+ padding-left: 0.85rem;
+}
diff --git a/_extra_zh/zh/code-highlight.js b/_extra_zh/zh/code-highlight.js
new file mode 100644
index 00000000..1e36669c
--- /dev/null
+++ b/_extra_zh/zh/code-highlight.js
@@ -0,0 +1,224 @@
+/* ── Code highlight — shared TIRx/Python tokenizer + code block builder ──
+ *
+ * Usage:
+ *
+ *
+ *
+ * API:
+ * highlightCode(text) — tokenize text, return HTML with tokens
+ * escapeHtml(text) — escape < > & for safe HTML insertion
+ *
+ * createCodeBlock(container, code, options)
+ * Renders highlighted code into container with optional regions + focus lines.
+ * options:
+ * blockDefs: [{ key, start, end, color }] — clickable regions
+ * focusLines: { key: [lineNos] } or { key: { lines, color } } — per-region line highlights
+ * onBlockClick: function(key) — callback on region click
+ *
+ * Auto-init:
+ * Elements with class "auto-hl" get their textContent highlighted on load.
+ * Combine with code-dark / code-light for theme:
+ * Tx.copy_async(Asmem, A[...])
+ *
+ * Panel structure (optional wrapper for header + rounded corners):
+ *
+ * Title
+ *
+ *
+ *
+ * Token classes (same as code-highlight.css):
+ * .kw — keyword .fn — callable .str — string
+ * .num — number .cmt — comment .op — operator
+ * .typ — type .dec — decorator
+ *
+ * Namespace convention: only the callable name gets .fn, NOT the prefix.
+ * Tx.copy_async → Tx.copy_async
+ * T.ptx.tcgen05.mma → T.ptx.tcgen05.mma
+ */
+
+(function (root) {
+ "use strict";
+
+ function escapeHtml(text) {
+ return text
+ .replace(/&/g, "&")
+ .replace(//g, ">");
+ }
+
+ var KEYWORDS = /^(def|with|for|in|if|else|elif|and|or|not|return|True|False|None|range|class|import|from|as|pass|break|continue|while|try|except|finally|raise|yield|lambda|assert|del|global|nonlocal)$/;
+
+ function classifyToken(tok) {
+ if (tok.startsWith("#")) return "cmt";
+ if (tok.startsWith("@")) return "dec";
+ if (tok.startsWith('"') || tok.startsWith("'")) return "str";
+ if (/^\d[\d.]*$/.test(tok)) return "num";
+ if (KEYWORDS.test(tok)) return "kw";
+ if (/^(?:Tx|T)(\.[A-Za-z_]\w*)+$/.test(tok)) return "fn";
+ if (/^[A-Za-z_]\w*$/.test(tok)) return "fn";
+ return "";
+ }
+
+ var TOKEN_RE = /("(?:[^"\\]|\\.)*"|'(?:[^'\\]|\\.)*'|#.*$|@[A-Za-z_][\w.]*|\b(?:def|with|for|in|if|else|elif|and|or|not|return|True|False|None|range|class|import|from|as|pass|break|continue|while|try|except|finally|raise|yield|lambda|assert|del|global|nonlocal)\b|\b\d[\d.]*\b|(?:Tx|T)(?:\.[A-Za-z_]\w*)+|\b[A-Za-z_]\w*(?=\())/gm;
+
+ function highlightCode(text) {
+ TOKEN_RE.lastIndex = 0;
+ var out = "";
+ var last = 0;
+ var m;
+ while ((m = TOKEN_RE.exec(text)) !== null) {
+ out += escapeHtml(text.slice(last, m.index));
+ var tok = m[0];
+ var cls = classifyToken(tok);
+ if (cls === "fn" && tok.indexOf(".") !== -1) {
+ // Namespace.callable — only wrap the last component
+ var lastDot = tok.lastIndexOf(".");
+ out += escapeHtml(tok.substring(0, lastDot + 1));
+ out += '' + escapeHtml(tok.substring(lastDot + 1)) + "";
+ } else if (cls) {
+ out += '' + escapeHtml(tok) + "";
+ } else {
+ out += escapeHtml(tok);
+ }
+ last = TOKEN_RE.lastIndex;
+ }
+ out += escapeHtml(text.slice(last));
+ return out;
+ }
+
+ // Auto-init: highlight elements with class "auto-hl"
+ if (typeof document !== "undefined") {
+ document.addEventListener("DOMContentLoaded", function () {
+ var els = document.querySelectorAll(".auto-hl");
+ for (var i = 0; i < els.length; i++) {
+ els[i].innerHTML = highlightCode(els[i].textContent);
+ }
+ });
+ }
+
+ // ── createCodeBlock: render code with optional regions + focus lines ──
+
+ function hexToRgb(hex) {
+ return [
+ parseInt(hex.slice(1, 3), 16),
+ parseInt(hex.slice(3, 5), 16),
+ parseInt(hex.slice(5, 7), 16)
+ ];
+ }
+
+ function createCodeBlock(container, code, options) {
+ options = options || {};
+ var blockDefs = options.blockDefs;
+ var focusLines = options.focusLines;
+ var onBlockClick = options.onBlockClick;
+ var lines = code.split("\n");
+
+ if (blockDefs) {
+ // Render code split into clickable regions
+ container.innerHTML = blockDefs.map(function (b) {
+ var lineHtml = lines.slice(b.start - 1, b.end).map(function (line, idx) {
+ var lineNo = b.start + idx;
+ var content = line.length ? highlightCode(line) : " ";
+ return '' + content + "";
+ }).join("");
+ var style = b.color ? ' style="border-left-color:' + b.color + '"' : "";
+ return '" + lineHtml + "";
+ }).join("");
+
+ // Color map for hover/active/focus
+ var colorMap = {};
+ blockDefs.forEach(function (b) { if (b.color) colorMap[b.key] = b.color; });
+
+ var regions = container.querySelectorAll(".cb-region");
+ var allLines = container.querySelectorAll(".cb-line");
+ var activeKey = null;
+
+ function clearFocus() {
+ for (var i = 0; i < allLines.length; i++) {
+ allLines[i].classList.remove("cb-focus");
+ allLines[i].style.background = "";
+ }
+ }
+
+ function applyFocus(key) {
+ if (!focusLines || !focusLines[key]) return;
+ var fl = focusLines[key];
+ var lns = Array.isArray(fl) ? fl : fl.lines;
+ var color = Array.isArray(fl) ? null : fl.color;
+ if (!color && colorMap[key]) {
+ var rgb = hexToRgb(colorMap[key]);
+ color = "rgba(" + rgb.join(",") + ",0.18)";
+ }
+ if (!color) color = "rgba(255,255,255,0.15)";
+ for (var i = 0; i < lns.length; i++) {
+ var el = container.querySelector('.cb-line[data-ln="' + lns[i] + '"]');
+ if (el) {
+ el.classList.add("cb-focus");
+ el.style.background = color;
+ }
+ }
+ }
+
+ for (var r = 0; r < regions.length; r++) {
+ (function (el) {
+ var regionKey = el.dataset.region;
+
+ // Hover with region color
+ if (colorMap[regionKey]) {
+ var rgb = hexToRgb(colorMap[regionKey]);
+ var hoverBg = "rgba(" + rgb.join(",") + ",0.15)";
+ el.addEventListener("mouseenter", function () {
+ if (!el.classList.contains("active")) el.style.background = hoverBg;
+ });
+ el.addEventListener("mouseleave", function () {
+ if (!el.classList.contains("active")) el.style.background = "";
+ });
+ }
+
+ // Click to activate
+ el.addEventListener("click", function () {
+ if (regionKey === activeKey) return;
+ activeKey = regionKey;
+ for (var j = 0; j < regions.length; j++) {
+ regions[j].classList.remove("active");
+ regions[j].style.background = "";
+ }
+ el.classList.add("active");
+ if (colorMap[regionKey]) {
+ var rgb2 = hexToRgb(colorMap[regionKey]);
+ el.style.background = "rgba(" + rgb2.join(",") + ",0.15)";
+ }
+ clearFocus();
+ applyFocus(regionKey);
+ if (onBlockClick) onBlockClick(regionKey);
+ });
+ })(regions[r]);
+ }
+ } else {
+ // Simple highlighted code, no regions
+ container.innerHTML = lines.map(function (line, idx) {
+ var content = line.length ? highlightCode(line) : " ";
+ return '' + content + "";
+ }).join("");
+ }
+ }
+
+ // ── renderAnnotatedCode: highlight with per-line CSS classes ──
+
+ function renderAnnotatedCode(container, code, lineClasses) {
+ lineClasses = lineClasses || {};
+ var lines = code.split("\n");
+ container.innerHTML = lines.map(function (line, idx) {
+ var content = line.length ? highlightCode(line) : " ";
+ var cls = lineClasses[idx + 1];
+ if (cls) return '' + content + "";
+ return "" + content + "";
+ }).join("");
+ }
+
+ // Export
+ root.highlightCode = highlightCode;
+ root.escapeHtml = escapeHtml;
+ root.createCodeBlock = createCodeBlock;
+ root.renderAnnotatedCode = renderAnnotatedCode;
+})(typeof window !== "undefined" ? window : this);
diff --git a/_extra_zh/zh/demo/barrier_intro.html b/_extra_zh/zh/demo/barrier_intro.html
new file mode 100644
index 00000000..9973f691
--- /dev/null
+++ b/_extra_zh/zh/demo/barrier_intro.html
@@ -0,0 +1,387 @@
+
+
+
+
+异步协作:Barrier
+
+
+
+
+
+
+异步协作:Barrier
+硬件单元完成后 signal mbarrier;warp group 在继续前等待
+
+
+
+
+ GMEM
A, B
+
+
+ TMA 加载
+
+ SMEM
+
+
+ tcgen05.mma
+
+ TMEM
累加器
+
+
+ tcgen05.ld
+
+ 寄存器
cast
+
+
+ 存储
+
+ SMEM
+
+
+ TMA 存储
+
+ GMEM
D
+
+
+
+
+
+
+
+
+
+
+
+
+
+ TMA
引擎
+
+
+ complete-tx
+
+
+ Tensor
Core
+
+
+
+
+
+
+
+
+
+
+
+
+ - mbarrier (mbar) — 用于协调线程和硬件单元之间的异步操作
+ - 不同 warp 可并行运行在操作的不同阶段
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/cta_cluster.html b/_extra_zh/zh/demo/cta_cluster.html
new file mode 100644
index 00000000..47e16dbb
--- /dev/null
+++ b/_extra_zh/zh/demo/cta_cluster.html
@@ -0,0 +1,113 @@
+
+
+
+
+
+2-CTA Cluster
+
+
+
+
+
+2-CTA Cluster:通过跨 CTA SMEM 读取实现协作 MMA
+点击任意部分查看说明;两个 CTA 在 cluster 内共享 stored-B 行切片(DSMEM)
+
+
+ CLUSTER — 跨 SM 的 2 个 CTA · DSMEM
+
+
+ CTA 0 · SM-0
+ Asmem(本 CTA)A 行 0–127
+ BsmemB stored 行 0–127
+ D[0:128, 0:256]
+
+
+ 跨 CTA
读取 ↔
+
+
+ CTA 1 · SM-1
+ Asmem(本 CTA)A 行 128–255
+ BsmemB stored 行 128–255
+ D[128:256, 0:256]
+
+
+ Cluster 输出:256 × 256 — 是单个 CTA 独立计算 128 × 128 的两倍
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/mbarrier_arrive_timeline.html b/_extra_zh/zh/demo/mbarrier_arrive_timeline.html
new file mode 100644
index 00000000..e2856809
--- /dev/null
+++ b/_extra_zh/zh/demo/mbarrier_arrive_timeline.html
@@ -0,0 +1,470 @@
+
+
+
+
+
+MBarrier arrive 时间线
+
+
+
+
+
+ 用于跨线程同步的 MBarrier
+
+
+
+
+
+
+
+ T0(生产者)
+
+
+
+
+
+
+
+ T1(消费者)
+
+
+
+
+ 时间线
+
+
+
+
+
+
+
+
+
+
+
+
+
+ 工作 / 活跃
+
+
+ Signal 操作
+
+
+ 阻塞 / 等待
+
+
+ Barrier 状态
+
+
+ 空闲
+
+
+ Signal 并改变 mbar 状态
+
+
+ init 后同步线程
+
+
+
+
+
+ 设置(T0)
+ mbarrier.init(mbar, 1) — 对所有运行只设置一次,设置后显式同步。
+
+
+ 生产者(T0)
+ mbarrier.arrive(mbar) 递减 pending_count;当它变为 0 时,phase 完成。
+
+
+ 消费者(T1)
+ mbarrier.try_wait(mbar, phase) 挂起 T1,直到 phase 完成。
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/mbarrier_mechanism.html b/_extra_zh/zh/demo/mbarrier_mechanism.html
new file mode 100644
index 00000000..16dbb513
--- /dev/null
+++ b/_extra_zh/zh/demo/mbarrier_mechanism.html
@@ -0,0 +1,220 @@
+
+
+
+
+
+MBarrier 机制
+
+
+
+
+
+ MBarrier:数据结构与 API
+
+
+
+ MBarrier 对象(64-bit,位于 shared memory)
+
+ phase0 或 1
+ pending_count剩余 arrival
+ expected_count预期总数
+ tx-count待完成字节
+
+
+ phase
+ 当前 phase(0 或 1)。完成后自动翻转。
+
+
+ pending_count
+ phase 完成前仍需要的 arrival 数量。
+
+
+ expected_count
+ 每个 phase 预期的 arrival 总数(由 init 设置)。
+
+
+ tx-count
+ 尚未完成的异步字节数。由 arrive.expect_tx 增加,由硬件在传输完成时减少。
+
+
+
+
+ Phase 完成条件
+
+ pending_count == 0
+ &&
+ tx-count ≤ 0
+
+
+ 当 所有预期线程都已 arrive 且 所有预期异步字节都已传输完成 时,一个 phase 完成。
+ 完成后,barrier 会自动重置:phase 翻转,pending_count 重置为 expected_count,tx-count 重置为 0。
+
+
+
+
+
+ 核心 API
+
+
+ 指令 角色 对 MBarrier 的影响
+
+
+
+ mbarrier.init 设置
+ 用预期 arrival 数初始化 barrier
+ phase=0, pending=count, tx=0
+
+
+ mbarrier.arrive 生产者
+ 线程向 barrier signal arrival
+ pending_count -= 1
+
+
+ mbarrier.arrive.expect_tx 生产者
+ Arrive 并声明预期异步传输字节数
+ pending -= 1, tx_count += txCount
+
+
+ tcgen05.commit 生产者
+ Tensor Core signal MMA 完成
+ 在 mbarrier 上执行 arrive::one
+
+
+ mbarrier.try_wait 消费者
+ 等待(挂起线程)直到 phase 完成
+ 阻塞直到 phase 完成
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/mbarrier_tcgen05_timeline.html b/_extra_zh/zh/demo/mbarrier_tcgen05_timeline.html
new file mode 100644
index 00000000..6889526f
--- /dev/null
+++ b/_extra_zh/zh/demo/mbarrier_tcgen05_timeline.html
@@ -0,0 +1,479 @@
+
+
+
+
+
+MBarrier tcgen05.commit 时间线
+
+
+
+
+
+ 用 MBarrier signal tcgen05 完成
+
+
+
+
+
+
+
+ T0(生产者)
+
+
+
+ Tensor Core
+
+
+
+
+
+
+
+ T1(消费者)
+
+
+
+
+ 时间线
+
+
+
+
+
+
+
+
+
+
+
+
+
+ 工作 / 活跃
+
+
+ Signal 操作
+
+
+ 派发到 Tensor Core
+
+
+ 阻塞 / 等待
+
+
+ Barrier 状态
+
+
+ 空闲
+
+
+ Signal 并改变 mbar 状态
+
+
+ init 后同步线程
+
+
+
+
+
+ 生产者(T0)
+ tcgen05.mma 发起异步矩阵乘。tcgen05.commit(mbar) 告诉 barrier 跟踪它;硬件完成后会执行 arrive::one。
+
+
+ Tensor Core
+ 执行 T0 排队的 tcgen05.mma 和 tcgen05.commit。
+
+
+ 消费者(T1)
+ mbarrier.try_wait(mbar, phase) 会挂起 T1,直到 phase 完成。结果随后可在 TMEM 中读取。
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/mbarrier_tma_timeline.html b/_extra_zh/zh/demo/mbarrier_tma_timeline.html
new file mode 100644
index 00000000..6ae2808a
--- /dev/null
+++ b/_extra_zh/zh/demo/mbarrier_tma_timeline.html
@@ -0,0 +1,490 @@
+
+
+
+
+
+MBarrier TMA 时间线
+
+
+
+
+
+ 用 MBarrier signal TMA 完成
+
+
+
+
+
+
+
+ T0(生产者)
+
+
+
+ TMA 引擎
+
+
+
+
+
+
+
+ T1(消费者)
+
+
+
+
+ 时间线
+
+
+
+
+
+
+
+
+
+
+
+
+
+ 工作 / 活跃
+
+
+ Signal 操作
+
+
+ 派发到 TMA
+
+
+ 阻塞 / 等待
+
+
+ Barrier 状态
+
+
+ 空闲
+
+
+ Signal 并改变 mbar 状态
+
+
+ init 后同步线程
+
+
+
+
+
+ 生产者(T0)
+ arrive.expect_tx(4096) 将 pending 递减到 0,并设置 tx_count。cp.async.bulk 触发 TMA 传输。
+
+
+ TMA 引擎
+ 执行从 T0 派发的 cp.async.bulk。完成时自动递减 tx_count。
+
+
+ 消费者(T1)
+ try_wait 会挂起,直到 pending==0 且 tx==0。硬件在传输完成时自动递减 tx_count。
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/phase_tracking.html b/_extra_zh/zh/demo/phase_tracking.html
new file mode 100644
index 00000000..94c82c34
--- /dev/null
+++ b/_extra_zh/zh/demo/phase_tracking.html
@@ -0,0 +1,201 @@
+
+
+
+
+
+Phase 跟踪
+
+
+
+
+
+ Phase 跟踪
+ Barrier 通过在 phase 0 ↔ 1 间翻转,在多次迭代之间复用
+
+
+
+
+
+
+
+ phase = 0迭代 0
+ phase = 1迭代 1
+ phase = 0迭代 2
+ phase = 1迭代 3
+ phase = 0迭代 4
+ phase = 1迭代 5
+ phase = 0迭代 6
+ phase = 1迭代 7
+
+
+
+
+ 工作方式
+
+ 每个 barrier 都有一个 phase bit(0 或 1)。所有预期 arrival 到达后,barrier 会翻转 phase 并重置,准备好服务下一次迭代;无需重新初始化。
+
+ 1 生产者 arrive 到 barrier(phase 0)
+ 2 消费者调用 try_wait(phase=0) — 解除阻塞
+ 3 Barrier 翻转到 phase 1,phase ^= 1
+
+
+
+
+
+ 为什么要跟踪 phase?
+
+ 在流水线 kernel 中,barrier 会在 K-loop 的每次迭代中复用。phase 跟踪解决的问题是:“这次 barrier 触发对应的是迭代 i,还是迭代 i-1?”
+
+ phase_tma = 0
+ for k in range(K):
+ try_wait(tma_bar, phase_tma)
+ phase_tma ^= 1 # 翻转
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/pipeline_arch.html b/_extra_zh/zh/demo/pipeline_arch.html
new file mode 100644
index 00000000..469773bc
--- /dev/null
+++ b/_extra_zh/zh/demo/pipeline_arch.html
@@ -0,0 +1,509 @@
+
+
+
+
+
+Blackwell 上的 GEMM 数据流水线
+
+
+
+
+
+ Blackwell 上的 GEMM 数据流水线
+ D = A × B — 点击一个操作,在架构图中查看对应的数据路径
+
+
+
+
+ GMEM
A, B
+
+
+
+ TMA 加载
+
+
+ SMEM
+
+
+
+ tcgen05.mma
+
+
+ TMEM
累加器
+
+
+
+ tcgen05.ld
+
+
+ 寄存器
cast
+
+
+
+ 存储
+
+
+ SMEM
+
+
+
+ TMA 存储
+
+
+ GMEM
D
+
+
+
+
+
+ Blackwell 架构
+
+
+
+
+
+
+
+ 流式多处理器 (SM)
+
+
+
+ Tensor Core (tcgen05)
+ 第 5 代 MMA
+
+
+
+ CUDA Core
+ FP/INT 单元
+
+
+
+
+
+ Shared Memory (SMEM)
+ 每个 SM 228 KB
+
+
+ Tensor Memory (TMEM)
+ 128 条 lane
+
+
+ 寄存器文件
+
+
+
+
+
+ TMA 引擎
+ 数据搬运器
+
+
+
+
+
+
+ SM ...
+ × N
+
+
+
+
+ 全局内存 (GMEM)
+
+
+
+
+
+
+
+
+
+ 点击流水线中的一个操作以高亮对应的数据路径。
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/sf_tmem.html b/_extra_zh/zh/demo/sf_tmem.html
new file mode 100644
index 00000000..48093a83
--- /dev/null
+++ b/_extra_zh/zh/demo/sf_tmem.html
@@ -0,0 +1,148 @@
+
+
+
+
+TMEM 中的 Scale Factor
+
+
+
+
+
+TMEM 中的 Scale Factor(nvfp4)
+M = 128,一个 MMA-K block(SF_K = 4) · warpx4 broadcast R[4 : 32@TLane]
+
+
+
+
+ 逻辑 SFA (M × SF_K)
+ 点击 cell = SFA[m, sfk] · 颜色 = m // 32 分组
+
+
+
+ TMEM (128 条 lane × 16 字节)
+ w0: m 0–31 · w1: m 32–63 · w2: m 64–95 · w3: m 96–127 · byte = sfk
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/sm_architecture.html b/_extra_zh/zh/demo/sm_architecture.html
new file mode 100644
index 00000000..77de85ca
--- /dev/null
+++ b/_extra_zh/zh/demo/sm_architecture.html
@@ -0,0 +1,397 @@
+
+
+
+
+
+Blackwell SM 架构
+
+
+
+
+
+ Blackwell SM 架构
+
+
+
+
+ 流式多处理器 (SM)
+
+
+
+
+ Tensor Core (tcgen05)
+ 第 5 代 MMA
+
+
+
+ CUDA Core
+ FP/INT 单元
+
+
+
+
+
+
+ Shared Memory (SMEM)
+ 每个 SM 228 KB
+
+
+ Tensor Memory (TMEM)
+ TMEM — 128 条 lane
+
+
+ 寄存器文件
+
+
+
+
+
+
+
+ TMA 引擎
+ 数据搬运器
+
+
+
+
+
+
+ SM ...
+ × N
+
+
+
+
+
+ 全局内存 (GMEM)
+
+
+
+
+
+
+
+
+
+ 点击硬件单元查看详情。实线箭头表示加载路径,虚线箭头表示存储路径。
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/swizzle_128B.html b/_extra_zh/zh/demo/swizzle_128B.html
new file mode 100644
index 00000000..e33b8886
--- /dev/null
+++ b/_extra_zh/zh/demo/swizzle_128B.html
@@ -0,0 +1,432 @@
+
+
+
+
+Swizzle 布局 — SWIZZLE_128B
+
+
+
+
+
+
+Swizzle 布局 — SWIZZLE_128B
+physical_sector = logical_sector XOR row | 8 行 × 32 个 bank,每个 sector 4 个 bank
+
+
+ 读取
+
+
+
+
+
+ 索引
+
+
+
+
+ 不使用 Swizzle
+
+ 32 个 Bank
+
+
+
+ SWIZZLE_128B
+
+ 32 个 Bank
+
+
+
+
+
+
+
+ 按 cycle 展示读取
+
+
+ 不使用 Swizzle
+ SWIZZLE_128B
+
+
+
+
+ 1 个元素 = 1 个 bank(4 字节)。一个 32 线程 warp 读取 32 个元素;bank conflict 会被串行化成额外 cycle。
+ 列 tile (8×4):8 行 × 4 列(1 个 sector)= 32 次读取。
+ 使用 swizzle:1 个 cycle;不使用:8 个 cycle(8 路冲突)。
+ 行 tile (1×32):1 行 × 32 列 = 32 次读取。
+ 始终是 1 个 cycle(同一行内 32 个 bank 全部不同)。
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/swizzle_8x8.html b/_extra_zh/zh/demo/swizzle_8x8.html
new file mode 100644
index 00000000..def1575e
--- /dev/null
+++ b/_extra_zh/zh/demo/swizzle_8x8.html
@@ -0,0 +1,384 @@
+
+
+
+
+Swizzle 布局
+
+
+
+
+
+
+Swizzle 布局 — 8×8
+mapped_col = logical_col XOR row | 8×8 矩阵,1 个元素 = 1 个 bank
+
+
+ 读取
+
+
+
+
+
+ 索引
+
+
+
+
+ 不使用 Swizzle
+ Bank 活跃情况
+
+
+ 使用 Swizzle(XOR)
+ Bank 活跃情况
+
+
+
+
+
+
+ 按 cycle 展示读取
+
+
+ 不使用 Swizzle
+ 使用 Swizzle(XOR)
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/swizzle_atom_general.html b/_extra_zh/zh/demo/swizzle_atom_general.html
new file mode 100644
index 00000000..fd2e7379
--- /dev/null
+++ b/_extra_zh/zh/demo/swizzle_atom_general.html
@@ -0,0 +1,472 @@
+
+
+
+
+Swizzle Atom — 全部格式
+
+
+
+
+
+
+Swizzle Atom — 全部格式
+swizzle atom 是内存中应用 swizzle 的基本连续区域
+
+
+
+ 格式
+
+
+
+
+
+
+
+
+
+
+
+
+ 读取
+
+
+
+
+
+ 索引
+
+
+ dtype
+
+
+
+
+
+
+
+
+
+
+
+ 不使用 Swizzle
+ Bank sector 活跃情况(全部 8 个)
+
+
+ 使用 Swizzle(XOR)
+ Bank sector 活跃情况(全部 8 个)
+
+
+
+
+
+关键:使用 swizzle 后,读取任意 8×16B 列都没有冲突!
+
+
+ 按 cycle 展示读取
+
+
+ 不使用 Swizzle
+ 使用 Swizzle(XOR)
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tcgen05_intro.html b/_extra_zh/zh/demo/tcgen05_intro.html
new file mode 100644
index 00000000..90692ea0
--- /dev/null
+++ b/_extra_zh/zh/demo/tcgen05_intro.html
@@ -0,0 +1,331 @@
+
+
+
+
+Tensor Core:tcgen05.mma
+
+
+
+
+
+
+tcgen05.mma
+
+
+ 转置 A
+
+
+
+
+
+ 转置 B
+
+
+
+
+
+
+
+ M(输出行)
+
+
+
+
+ N(输出列)
+
+
+
+
+
+
+
+ K
+
+
+
+
+
+
+ 行分组
+ 列分组
+
+
+
+
+ A (SMEM)
+
+
+
+ ×
+
+ B (SMEM)
+
+
+
+
+
+ Tensor Core
+ tcgen05.mma
+
+ →
+
+
+ C (TMEM)
+
+
+ 128 条 lane(M)
+
+
+ 列(N)— 由 tmem.alloc 分配
+
+
+
+
+
+
+
+
+
+
+
+
+ SMEM 描述符(A, B)
+ smem_desc = base_addr | leading_byte_offset | stride_byte_offset | start_addr_off
+
+
+ tcgen05.mma PTX
+ tcgen05.mma.cta_group::1.kind::f16 taddr, a_desc, b_desc, idesc, enable;
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/thread_hierarchy.html b/_extra_zh/zh/demo/thread_hierarchy.html
new file mode 100644
index 00000000..8c49996f
--- /dev/null
+++ b/_extra_zh/zh/demo/thread_hierarchy.html
@@ -0,0 +1,112 @@
+
+
+
+
+
+Blackwell 线程层级
+
+
+
+
+
+Blackwell 线程层级
+点击一个层级:thread → warp → warpgroup → CTA → cluster → grid
+
+
+
+ GRID — 一次 kernel launch
+
+ CLUSTER — 跨 SM 的 CTA(DSMEM)
+
+
+ CTA — 线程块 · 一个 SM
+
+ WARPGROUP — 4 个 warp · 128 个线程
+
+
+ WARP — 32 个线程(SIMT)
+
+
+ warp 1
+ warp 2
+ warp 3
+
+
+
+ CTA 1
…
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/thread_register.html b/_extra_zh/zh/demo/thread_register.html
new file mode 100644
index 00000000..99029747
--- /dev/null
+++ b/_extra_zh/zh/demo/thread_register.html
@@ -0,0 +1,252 @@
+
+
+
+
+线程 + 寄存器布局
+
+
+
+
+
+
+Layout: S[(8, 4, 2) : (4@laneid, 1@laneid, 1@reg)]
+点击 cell 查看分配关系
+
+
+
+ 逻辑 8×8 矩阵
+
+
+
+
+ 线程分配
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tile_distributed.html b/_extra_zh/zh/demo/tile_distributed.html
new file mode 100644
index 00000000..f0c45819
--- /dev/null
+++ b/_extra_zh/zh/demo/tile_distributed.html
@@ -0,0 +1,480 @@
+
+
+
+
+分布式布局
+
+
+
+
+
+
+分布式布局 — 2×2 GPU Mesh 上的 8×8 矩阵
+完全 shard 到 4 个 GPU | 点击 cell 查看放置位置
+
+
+ 布局
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+
+ GPU Mesh(2×2)
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tiled_layout.html b/_extra_zh/zh/demo/tiled_layout.html
new file mode 100644
index 00000000..fbe7c961
--- /dev/null
+++ b/_extra_zh/zh/demo/tiled_layout.html
@@ -0,0 +1,278 @@
+
+
+
+
+Tiled 布局
+
+
+
+
+
+
+S[(4, 2, 2, 4) : (16, 4, 8, 1)]
+逻辑视图 vs 物理内存 | 点击 cell 查看映射
+
+
+
+ 逻辑矩阵
+
+
+
+
+
+
+ 物理内存(2×4 tiled)
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tiling_constraint.html b/_extra_zh/zh/demo/tiling_constraint.html
new file mode 100644
index 00000000..16480f68
--- /dev/null
+++ b/_extra_zh/zh/demo/tiling_constraint.html
@@ -0,0 +1,343 @@
+
+
+
+
+Swizzle 带来的 tiling 约束
+
+
+
+
+
+
+Swizzle 带来的 tiling 约束
+在 16×16 矩阵上使用 128B swizzle — 切成 16×8 分组即可消除 bank conflict
+
+
+ 是否 tile
+
+
+
+
+
+ 行范围
+
+
+
+
+
+
+
+ 列
+
+
+
+
+ 原始 16×16 矩阵
+
+
+
+
+
+
+ 为 128B Swizzle 进行 tiling(两个 16×8 分组)
+
+
+
+
+
+ dtype
+
+
+
+
+
+
+
+
+
+ 硬件会基于物理偏移计算 swizzle_col,等价于假设每行只有 8 个元素。若每行有 16 个元素,则每个原始行在 swizzle 计算中会被当成 2 行处理。
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tirx_dispatch.html b/_extra_zh/zh/demo/tirx_dispatch.html
new file mode 100644
index 00000000..0470045b
--- /dev/null
+++ b/_extra_zh/zh/demo/tirx_dispatch.html
@@ -0,0 +1,159 @@
+
+
+
+
+
+TIRx:Scope、Layout、Dispatch
+
+
+
+
+
+TIRx:Scope、Layout、Dispatch
+点击一个设计元素,查看它控制 single-MMA GEMM 中的哪些代码行
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tma_3d.html b/_extra_zh/zh/demo/tma_3d.html
new file mode 100644
index 00000000..e118e70b
--- /dev/null
+++ b/_extra_zh/zh/demo/tma_3d.html
@@ -0,0 +1,353 @@
+
+
+
+
+用 3D TMA 做 tiling 与 swizzle
+
+
+
+
+
+
+用 3D TMA 做 tiling 与 swizzle
+使用 3D TMA 将数据复制到 tiled shared memory
+
+
+ Swizzle
+
+
+
+
+
+ 列偏移
+
+
+
+
+
+
+ Global Memory
+ 16×256 fp16(连续布局,row = 512B)
+
+
+
+
+ →
+ 3D TMA
+
+
+
+
+ Shared Memory(tiled & swizzle 后)
+
+
+
+
+
+
+
+
+
+
+
+
+
diff --git a/_extra_zh/zh/demo/tma_intro.html b/_extra_zh/zh/demo/tma_intro.html
new file mode 100644
index 00000000..65a79ed4
--- /dev/null
+++ b/_extra_zh/zh/demo/tma_intro.html
@@ -0,0 +1,358 @@
+
+
+
+
+TMA:Tensor Memory Accelerator
+
+
+
+
+
+
+TMA:从 Global Memory 到 Shared Memory
+硬件 2D copy,可选 swizzle — 一条指令,无需线程搬运
+
+
+ Swizzle
+
+
+
+
+
+ 行偏移
+ 列偏移
+
+
+
+
+
+
+ Global Memory
+ 16×128 fp16(每个 cell = 16B = 1×8 fp16)
+
+
+
+
+
+ TMA 引擎
+ 数据搬运器
+
+ →
+ cp.async.bulk.tensor.2d
+ SWIZZLE_128B
+
+
+
+ Shared Memory
+ 8×8 个 sector(swizzle 后)
+
+
+
+
+
+
+
+
+
+
+
+ - ▸ 单线程派发(非阻塞)
+ - ▸ 自动 swizzle — 数据到达时天然避免 bank conflict
+ - ▸ 双向:load(GMEM→SMEM)和 store(SMEM→GMEM)
+
+
+
+
+
diff --git a/_extra_zh/zh/viz-base.css b/_extra_zh/zh/viz-base.css
new file mode 100644
index 00000000..314d742d
--- /dev/null
+++ b/_extra_zh/zh/viz-base.css
@@ -0,0 +1,90 @@
+/* Shared base styles for all viz HTMLs
+ * Colors: --bg gray, --accent blue, --surface white
+ * Fonts: Inter (body), SF Mono (code/notation)
+ */
+
+:root {
+ --bg:#fff; --surface:#fff; --border:#dfe1e6; --text:#222; --dim:#888; --accent:#3b82f6;
+
+ /* ── Group palette (tiles, lanes, banks, sectors) ─────── */
+ --color-group-0: #5b9bd5;
+ --color-group-1: #ed9a3c;
+ --color-group-2: #d95555;
+ --color-group-3: #45b5a5;
+ --color-group-4: #5fb85f;
+ --color-group-5: #e0b828;
+ --color-group-6: #9d6eb5;
+ --color-group-7: #e87888;
+
+ /* ── Interaction ──────────────────────────────────────── */
+ --color-hover-bg: #dbeafe;
+ --color-hover-text: #1e40af;
+
+ /* ── Status ───────────────────────────────────────────── */
+ --color-good: #2e7d32;
+ --color-bad: #c62828;
+
+ /* ── Boundaries & neutrals ────────────────────────────── */
+ --color-boundary: #8C1515;
+ --color-cell-bg: #f0f1f3;
+ --color-cell-bg-alt: #e8eaed;
+}
+* { box-sizing:border-box; margin:0; padding:0; }
+body { background:var(--bg); color:var(--text); font-family:'Inter','SF Pro','Segoe UI',system-ui,sans-serif; padding:24px; }
+body.figure h1, body.figure .sub { display:none; }
+body.notitle h1, body.notitle .sub { display:none; }
+h1 { text-align:center; font-size:21px; font-weight:700; margin-bottom:2px; }
+.sub { text-align:center; color:var(--dim); font-size:14px; margin-bottom:20px; font-family:'SF Mono','Fira Code',monospace; }
+
+/* Controls */
+.controls { display:flex; gap:18px; justify-content:center; align-items:center; margin-bottom:10px; flex-wrap:wrap; }
+.lbl { font-size:13px; color:var(--dim); margin-right:4px; font-weight:600; }
+.bg { display:inline-flex; gap:2px; }
+.btn {
+ padding:5px 11px; border:1px solid var(--border); background:var(--surface); color:var(--text);
+ cursor:pointer; border-radius:5px; font-size:13px; font-family:inherit; font-weight:500; transition:all .12s;
+}
+.btn:hover { background:#eef0f4; }
+.btn.on { background:var(--accent); border-color:var(--accent); color:#fff; }
+.btn.on:hover { background:#2563eb; }
+
+/* Side-by-side panels */
+.panels { display:grid; grid-template-columns:1fr 1fr; gap:20px; max-width:1100px; margin:0 auto; position:relative; }
+.panel { background:var(--surface); border-radius:10px; padding:16px 14px; border:1px solid var(--border);
+ box-shadow:0 1px 3px rgba(0,0,0,.06); }
+.panel h2 { text-align:center; font-size:14px; font-weight:700; margin-bottom:2px; }
+.panel .nota { text-align:center; font-size:11px; color:var(--accent); font-family:'SF Mono','Fira Code',monospace;
+ margin-bottom:8px; font-weight:600; }
+
+/* Grid cells */
+.grid { display:grid; gap:3px; }
+.hdr { font-size:10px; color:var(--dim); text-align:center; padding:2px 0; font-weight:600; }
+.rl { font-size:10px; color:var(--dim); display:flex; align-items:center; justify-content:flex-end; padding-right:3px; font-weight:600; }
+.cell {
+ aspect-ratio:1; border-radius:5px; display:flex; flex-direction:column; align-items:center; justify-content:center;
+ font-size:13px; font-weight:700; border:2.5px solid transparent; transition:all .18s;
+ min-width:0; cursor:pointer; line-height:1.15;
+}
+.cell.hov { border-color:#222; border-width:3px; z-index:2; box-shadow:0 0 8px rgba(59,130,246,.4); }
+.cell.dm { opacity:.25; }
+
+/* Arrow SVG overlay */
+.arrow-svg { position:absolute; top:0; left:0; width:100%; height:100%; pointer-events:none; z-index:10; overflow:visible; }
+.arrow-svg path { fill:none; stroke:#222; stroke-width:1.5; }
+.arrow-svg text { font-size:11px; font-weight:600; fill:#222; font-family:inherit; }
+
+/* Formula bar */
+.formula-bar { max-width:1100px; margin:18px auto 0; background:var(--surface); border-radius:10px;
+ padding:12px 18px; border:1px solid var(--border); box-shadow:0 1px 3px rgba(0,0,0,.06); }
+.formula-bar .ftitle { font-size:13px; font-weight:700; margin-bottom:6px; }
+.formula-bar .fcontent { font-size:13px; font-family:'SF Mono','Fira Code',monospace; color:var(--accent); line-height:1.6; }
+
+/* Legend */
+.leg { display:flex; flex-direction:column; align-items:center; gap:4px; margin-top:14px; }
+.leg-row { display:flex; gap:12px; justify-content:center; flex-wrap:wrap; }
+.leg-group { display:flex; gap:10px; align-items:center; }
+.leg-sep { width:1px; height:16px; background:var(--border); }
+.li { display:flex; align-items:center; gap:4px; font-size:12px; color:var(--dim); }
+.swtch { width:14px; height:14px; border-radius:3px; }
+
+@media(max-width:700px) { .panels { grid-template-columns:1fr; } }
diff --git a/_extra_zh/zh/viz-base.js b/_extra_zh/zh/viz-base.js
new file mode 100644
index 00000000..ce3e1a3f
--- /dev/null
+++ b/_extra_zh/zh/viz-base.js
@@ -0,0 +1,66 @@
+// Shared behavior for all viz HTMLs
+document.addEventListener('DOMContentLoaded', function() {
+ var p = new URLSearchParams(location.search);
+ if (p.has('notitle')) document.body.classList.add('notitle');
+
+ // Directly opened demo pages are copied outside Sphinx templates, so they do
+ // not receive html_js_files. Load the same language switch used by book pages.
+ if (window.parent === window && !p.has('notitle')) {
+ var langSwitchPath = window.location.pathname.indexOf('/zh/demo/') >= 0
+ ? '../../_static/lang-switch.js'
+ : '../_static/lang-switch.js';
+ var s = document.createElement('script');
+ s.src = new URL(langSwitchPath, window.location.href).href;
+ document.head.appendChild(s);
+ }
+
+ // Forward arrow keys to parent (reveal.js) when embedded
+ if (window.parent !== window) {
+ document.addEventListener('keydown', function(e) {
+ if ([37, 38, 39, 40, 27, 32].indexOf(e.keyCode) !== -1) {
+ // Left, Up, Right, Down, Escape, Space
+ window.parent.postMessage({ type: 'revealKey', keyCode: e.keyCode }, '*');
+ }
+ });
+ }
+});
+
+// Auto-height: when embedded in the book (demo-embed.js), the demo measures its
+// OWN content height and posts it to the parent, which sizes the iframe to fit so
+// there is never an inner scrollbar. This is push-based on purpose — the demo
+// catches its own DOM changes (a click that appends rows, expands a panel, …),
+// which a parent watching the iframe's from outside can miss. Measuring
+// body.scrollHeight (not documentElement, which is floored to the viewport) lets
+// the reported height grow AND shrink with the content.
+(function () {
+ if (window.parent === window) return; // only when embedded
+ var lastH = 0;
+ function report() {
+ var b = document.body, de = document.documentElement;
+ var h = (b ? b.scrollHeight : 0) || (de ? de.scrollHeight : 0) || 0;
+ if (h && Math.abs(h - lastH) > 1) {
+ lastH = h;
+ window.parent.postMessage({ type: 'demoHeight', height: h }, '*');
+ }
+ }
+ var scheduled = false;
+ function schedule() {
+ if (scheduled) return;
+ scheduled = true;
+ requestAnimationFrame(function () { scheduled = false; report(); });
+ }
+ // documentElement exists even while we are still in , so observers can be
+ // attached immediately; the first read happens in the rAF after layout.
+ try { new ResizeObserver(schedule).observe(document.documentElement); } catch (e) {}
+ try {
+ new MutationObserver(schedule).observe(document.documentElement, {
+ subtree: true, childList: true, attributes: true, characterData: true
+ });
+ } catch (e) {}
+ document.addEventListener('DOMContentLoaded', schedule);
+ window.addEventListener('load', schedule);
+ // Clicks often trigger async content changes; re-measure right after.
+ window.addEventListener('click', function () { setTimeout(schedule, 0); }, true);
+ // Catch late settling (fonts, deferred render).
+ [100, 300, 600, 1200].forEach(function (t) { setTimeout(schedule, t); });
+})();
diff --git a/_templates/sbt-sidebar-nav.html b/_templates/sbt-sidebar-nav.html
new file mode 100644
index 00000000..1b39afbc
--- /dev/null
+++ b/_templates/sbt-sidebar-nav.html
@@ -0,0 +1,32 @@
+
diff --git a/conf.py b/conf.py
index 0a03f5a5..1315caf4 100644
--- a/conf.py
+++ b/conf.py
@@ -44,12 +44,13 @@
html_title = project
html_logo = "static/mlc-logo-with-text-landscape.svg"
html_favicon = "static/mlc-favicon.ico"
+templates_path = ["_templates"]
html_static_path = ["static"]
# Interactive slide demos (self-contained HTML+CSS+JS) copied verbatim into the
# site root, then embedded via