…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/code-highlight.js b/_extra/code-highlight.js
deleted file mode 100644
index 1e36669c..00000000
--- a/_extra/code-highlight.js
+++ /dev/null
@@ -1,224 +0,0 @@
-/* ── 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/demo/barrier_intro.html b/_extra/demo/barrier_intro.html
deleted file mode 100644
index a5d1d671..00000000
--- a/_extra/demo/barrier_intro.html
+++ /dev/null
@@ -1,387 +0,0 @@
-
-
-
-
-Async Coordination: Barriers
-
-
-
-
-
-
-Async Coordination: Barriers
-HW units signal mbarrier on completion — warp groups wait before proceeding
-
-
-
-
- GMEM
A, B
-
-
- TMA load
-
- SMEM
-
-
- tcgen05.mma
-
- TMEM
accum
-
-
- tcgen05.ld
-
- Reg
cast
-
-
- store
-
- SMEM
-
-
- TMA store
-
- GMEM
D
-
-
-
-
-
-
-
-
-
-
-
-
-
- TMA
Engine
-
-
- complete-tx
-
-
- Tensor
Core
-
-
-
-
-
-
-
-
-
-
-
-
- - mbarrier (mbar) — used to coordinate asynchronous operations across threads and hardware units
- - Different warps run in parallel on different stages of operations
-
-
-
-
-
diff --git a/_extra/demo/mbarrier_arrive_timeline.html b/_extra/demo/mbarrier_arrive_timeline.html
deleted file mode 100644
index 30be850e..00000000
--- a/_extra/demo/mbarrier_arrive_timeline.html
+++ /dev/null
@@ -1,470 +0,0 @@
-
-
-
-
-
-MBarrier Arrive Timeline
-
-
-
-
-
- MBarrier for Cross-Thread Synchronization
-
-
-
-
-
-
-
- T0 (producer)
-
-
-
-
-
-
-
- T1 (consumer)
-
-
-
-
- Timeline
-
-
-
-
-
-
-
-
-
-
-
-
-
- Working / active
-
-
- Signaling operation
-
-
- Blocked / waiting
-
-
- Barrier state
-
-
- Idle
-
-
- Signaling to change mbar state
-
-
- Sync threads after init
-
-
-
-
-
- Setup (T0)
- mbarrier.init(mbar, 1) — setup once across all runs, explicit sync after setup.
-
-
- Producer (T0)
- mbarrier.arrive(mbar) decrements pending_count. When it hits 0, phase completes.
-
-
- Consumer (T1)
- mbarrier.try_wait(mbar, phase) suspends T1 until phase completes.
-
-
-
-
-
-
diff --git a/_extra/demo/mbarrier_tcgen05_timeline.html b/_extra/demo/mbarrier_tcgen05_timeline.html
deleted file mode 100644
index 2e16f24e..00000000
--- a/_extra/demo/mbarrier_tcgen05_timeline.html
+++ /dev/null
@@ -1,479 +0,0 @@
-
-
-
-
-
-MBarrier tcgen05.commit Timeline
-
-
-
-
-
- MBarrier to Signal tcgen05 Completion
-
-
-
-
-
-
-
- T0 (producer)
-
-
-
- Tensor Core
-
-
-
-
-
-
-
- T1 (consumer)
-
-
-
-
- Timeline
-
-
-
-
-
-
-
-
-
-
-
-
-
- Working / active
-
-
- Signaling operation
-
-
- Dispatch to TensorCore
-
-
- Blocked / waiting
-
-
- Barrier state
-
-
- Idle
-
-
- Signaling to change mbar state
-
-
- Sync threads after init
-
-
-
-
-
- Producer (T0)
- tcgen05.mma issues async matrix multiply. tcgen05.commit(mbar) tells barrier to track it; HW will arrive::one when done.
-
-
- Tensor Core
- Executes tcgen05.mma and tcgen05.commit queued from T0.
-
-
- Consumer (T1)
- mbarrier.try_wait(mbar, phase) suspends T1 until phase completes. Result available in TMEM.
-
-
-
-
-
-
diff --git a/conf.py b/conf.py
index 5b61b998..92052dd9 100644
--- a/conf.py
+++ b/conf.py
@@ -34,9 +34,6 @@
"zh",
"_*.md",
"**/_*.md",
- "setup.py",
- "tirx_tutorial",
- "references.bib",
"img/scripts",
".git",
".github",
diff --git a/img/scripts/README.md b/img/scripts/README.md
index 7e7eadbc..6b246624 100644
--- a/img/scripts/README.md
+++ b/img/scripts/README.md
@@ -1,10 +1,9 @@
# Diagram Generation Scripts
-These scripts generate current and legacy tutorial diagrams. Run from this directory:
+These scripts generate the tutorial diagrams. Run from this directory:
```bash
cd img/scripts
-python gen_cross_warp_reduce.py # -> ../cross_warp_reduce.png
python gen_flash_attention_barrier_flow.py # -> ../flash_attention_main_handoff.png, ../flash_attention_softmax_correction.png
python gen_flash_attention_pipeline.py # -> ../flash_attention_pipeline_v2.png
python gen_gemm_perf.py # -> ../gemm_perf.png
@@ -13,8 +12,6 @@ python gen_memory_dataflow.py # -> ../memory_dataflow.png
python gen_mma_layouts.py # -> ../mma_cg1_m128.svg, ../mma_cg1_m64.svg, ../mma_cg2_m256.svg, ../mma_cg2_m128.svg, ../mma_block_scaled.svg
python gen_roofline.py # -> ../roofline.png
python gen_sf_scale_vec.py # -> ../sf_scale_vec.svg
-python gen_sf_tmem.py # -> ../sf_tmem.svg
-python gen_shuffle_reduce.py # -> ../shuffle_reduce.png
python gen_smem_descriptor.py # -> ../wgmma_descriptor_kmajor.svg
python gen_swizzle_conflict.py # -> ../swizzle_conflict.svg
python gen_tcgen05_ldst.py # -> ../tcgen05_ldst.svg
@@ -25,6 +22,4 @@ python gen_warp_specialization_timeline.py # -> ../warp_specialization_timelin
Requires: `matplotlib`, `numpy`
-The images referenced by the current tutorial are checked into `img/`. Some scripts are kept only
-for reproducibility of older or optional diagrams, so their outputs may not be checked in until
-they are needed again.
+The images referenced by the tutorial are checked into `img/`.
diff --git a/img/scripts/gen_cross_warp_reduce.py b/img/scripts/gen_cross_warp_reduce.py
deleted file mode 100644
index 79f394e1..00000000
--- a/img/scripts/gen_cross_warp_reduce.py
+++ /dev/null
@@ -1,90 +0,0 @@
-"""Generate Cross-Warp Reduction diagram (chapter_rmsnorm)."""
-import matplotlib
-matplotlib.use('Agg')
-import matplotlib.pyplot as plt
-import matplotlib.patches as mpatches
-import numpy as np
-
-fig, ax = plt.subplots(figsize=(18, 8))
-ax.axis('off')
-
-N_WARPS = 4
-warp_colors = ['#4a90d9', '#e67e22', '#27ae60', '#8e44ad']
-warp_sums = [120, 85, 200, 95]
-grand_total = sum(warp_sums)
-y_mid = (N_WARPS - 1) * 1.5 / 2
-
-# Step 1: Each warp writes to SMEM
-ax.text(0.5, -0.7, "Step 1: Write to SMEM", ha='center', fontsize=11, fontweight='bold')
-for w in range(N_WARPS):
- y = w * 1.5
- ax.text(0, y, f"Warp {w}\nsum={warp_sums[w]}", ha='center', va='center', fontsize=10,
- bbox=dict(boxstyle='round,pad=0.4', fc=warp_colors[w], ec='black', lw=1.5, alpha=0.3))
- ax.annotate("", xy=(1.5, y), xytext=(0.75, y),
- arrowprops=dict(arrowstyle='->', color=warp_colors[w], lw=2))
-
-# SMEM boxes
-x_smem = 2.0
-ax.text(x_smem, -0.7, "SMEM", ha='center', fontsize=11, fontweight='bold', color='#555')
-for w in range(N_WARPS):
- y = w * 1.5
- ax.text(x_smem, y, f"smem[{w}]={warp_sums[w]}", ha='center', va='center', fontsize=9,
- bbox=dict(boxstyle='round,pad=0.3', fc='#fff3cd', ec='#ffc107', lw=1.5))
-
-# Step 2: Barrier
-x_bar1 = 4.0
-ax.text(x_bar1, -0.7, "Step 2", ha='center', fontsize=11, fontweight='bold')
-ax.text(x_bar1, y_mid, "bar.sync\n────────\nall warps\nwait here", ha='center', va='center', fontsize=10,
- bbox=dict(boxstyle='round,pad=0.5', fc='#ffcccc', ec='#e74c3c', lw=2))
-
-# Step 3: Warp 0 reads all from SMEM
-x_read = 6.5
-ax.text(x_read, -0.7, "Step 3: Warp 0 reads", ha='center', fontsize=11, fontweight='bold')
-for w in range(N_WARPS):
- y = w * 1.5
- ax.text(x_read, y, f"t{w}: {warp_sums[w]}", ha='center', va='center', fontsize=10,
- bbox=dict(boxstyle='round,pad=0.3', fc='#e8f4fd', ec='#4a90d9', lw=1.5))
- ax.annotate("", xy=(x_read - 0.55, y), xytext=(x_smem + 0.7, y),
- arrowprops=dict(arrowstyle='->', color='#4a90d9', lw=1.5))
-ax.text(x_read, N_WARPS * 1.5, "Warp 0 only", ha='center', fontsize=9,
- style='italic', color='#666')
-
-# Step 4: Shuffle XOR reduce
-x_shuf = 9.5
-ax.text(x_shuf, -0.7, "Step 4: Shuffle reduce", ha='center', fontsize=11, fontweight='bold')
-ax.text(x_shuf, y_mid, f"Shuffle XOR\n(same pattern\nas before)\n\n→ {grand_total}",
- ha='center', va='center', fontsize=10,
- bbox=dict(boxstyle='round,pad=0.6', fc='#d5f5e3', ec='#27ae60', lw=2))
-ax.annotate("", xy=(x_shuf - 0.9, y_mid), xytext=(x_read + 0.55, y_mid),
- arrowprops=dict(arrowstyle='->', color='#4a90d9', lw=2))
-
-# Step 5: Write grand total to smem[0]
-x_write = 12.5
-ax.text(x_write, -0.7, "Step 5: Write total", ha='center', fontsize=11, fontweight='bold')
-ax.text(x_write, y_mid, f"smem[0]\n= {grand_total}", ha='center', va='center', fontsize=11,
- fontweight='bold',
- bbox=dict(boxstyle='round,pad=0.5', fc='#fff3cd', ec='#27ae60', lw=2))
-ax.annotate("", xy=(x_write - 0.6, y_mid), xytext=(x_shuf + 0.9, y_mid),
- arrowprops=dict(arrowstyle='->', color='#27ae60', lw=2))
-
-# Step 6: Barrier + all warps read
-x_final = 15.5
-ax.text(x_final, -0.7, "Step 6: Barrier + Read", ha='center', fontsize=11, fontweight='bold')
-for w in range(N_WARPS):
- y = w * 1.5
- ax.text(x_final, y, f"Warp {w}\n= {grand_total}", ha='center', va='center', fontsize=10,
- fontweight='bold',
- bbox=dict(boxstyle='round,pad=0.4', fc='#d5f5e3', ec='#27ae60', lw=1.5))
- ax.annotate("", xy=(x_final - 0.7, y), xytext=(x_write + 0.6, y_mid),
- arrowprops=dict(arrowstyle='->', color='#27ae60', lw=1.2,
- connectionstyle=f'arc3,rad={0.1 * (w - 1.5)}'))
-
-ax.set_title("Cross-Warp Reduction: Gather → Reduce → Broadcast",
- fontsize=14, fontweight='bold', pad=20)
-ax.set_xlim(-1.5, 17)
-ax.set_ylim(-1.3, N_WARPS * 1.5 + 0.5)
-ax.invert_yaxis()
-plt.tight_layout()
-plt.savefig('../cross_warp_reduce.png', dpi=150, bbox_inches='tight')
-print('Saved cross_warp_reduce.png')
-plt.close()
diff --git a/img/scripts/gen_sf_tmem.py b/img/scripts/gen_sf_tmem.py
deleted file mode 100644
index 26d88c5a..00000000
--- a/img/scripts/gen_sf_tmem.py
+++ /dev/null
@@ -1,107 +0,0 @@
-"""Block-scaled MMA scale factors (SFA, SFB) in TMEM. Ground truth: nymph-rust +
-tvm sf_tmem_layout (backend/cuda/.../gemm_async/tcgen05.py): rows must be a
-multiple of 32; M = rows // 32; epc = 4 (four 8-bit SFs per 32-bit TMEM column);
-the atom is one 32-row chunk with R[4 : 32@TLane] (a broadcast to 4 warps).
-
-Two distinct mappings, drawn as two panels:
- (1) Packing — 128 M-rows occupy only 32 TMEM lanes (TLane = m % 32; the m // 32
- group runs along TCol).
- (2) Replication — those 32 stored lanes are broadcast to 4 warps (R[4 : 32@TLane])
- to all 128 lanes of the reading warpgroup: lane l reads TLane (l mod 32).
-Outputs SVG (and a PNG for inspection).
-"""
-import matplotlib
-matplotlib.use("Agg")
-import matplotlib.pyplot as plt
-from matplotlib.patches import Rectangle, FancyArrowPatch
-
-from pathlib import Path; OUT = str(Path(__file__).resolve().parent.parent) # the repo img/ dir
-TXT = "#1f2937"
-PURPLE = ["#7c3aed", "#8b5cf6", "#a78bfa", "#c4b5fd"]
-
-fig, ax = plt.subplots(figsize=(11.6, 5.7))
-fig.patch.set_facecolor("white")
-ax.set_xlim(0, 100)
-ax.set_ylim(0, 100)
-ax.axis("off")
-ax.text(50, 98.5, "Scale factors in TMEM — packed into 32 lanes, then broadcast to 4 warps",
- ha="center", va="top", fontsize=12.8, fontweight="bold", color=TXT)
-
-# ----------------------------------------------------------------------------
-# Panel 1 — packing: 128 M-rows -> 32 TMEM lanes
-# ----------------------------------------------------------------------------
-ax.text(22, 90, "① Packed: 128 M-rows → 32 TMEM lanes", ha="center", fontsize=9.6,
- fontweight="bold", color=TXT)
-ax.text(22, 86.3, "TLane = m % 32 m // 32 → TCol", ha="center", fontsize=7.6,
- color=TXT, style="italic")
-
-X0, CW, YT, RH = 12, 7.4, 79, 8.0
-rows = [0, 1, 2, None, 31]
-for gi in range(4): # m // 32 group -> TCol
- cx = X0 + gi * CW
- ax.text(cx + CW / 2, YT + 1.3, str(gi), ha="center", va="bottom", fontsize=7,
- fontweight="bold", color=TXT)
- for ri, lane in enumerate(rows):
- y = YT - (ri + 1) * RH
- if lane is None:
- ax.text(cx + CW / 2, y + RH / 2, "⋮", ha="center", va="center", fontsize=11, color=TXT)
- continue
- ax.add_patch(Rectangle((cx, y), CW, RH, facecolor=PURPLE[gi], edgecolor="white",
- linewidth=1.3, alpha=0.92))
- ax.text(cx + CW / 2, y + RH / 2, f"r{gi * 32 + lane}", ha="center", va="center",
- color="white", fontsize=6.4, fontweight="bold")
-for ri, lane in enumerate(rows): # left lane labels
- y = YT - (ri + 1) * RH
- lab = "⋮" if lane is None else f"TLane {lane}"
- ax.text(X0 - 1.3, y + RH / 2, lab, ha="right", va="center", fontsize=7, fontweight="bold", color=TXT)
-ax.text(X0 + 2 * CW, YT + 4.6, "TCol → (m // 32 group, then K)", ha="center", fontsize=7.3, color=TXT)
-ax.text(22, 16, "Only 32 lanes hold all 128 M-rows.", ha="center", fontsize=7.6, color=TXT, style="italic")
-
-# ----------------------------------------------------------------------------
-# Bridge arrow — broadcast to 4 warps
-# ----------------------------------------------------------------------------
-ax.annotate("", xy=(50.5, 52), xytext=(43.5, 52),
- arrowprops=dict(arrowstyle="-|>", color="#7c3aed", lw=2.4))
-ax.text(47, 56, "broadcast to\n4 warps", ha="center", va="bottom", fontsize=7.6,
- fontweight="bold", color="#7c3aed")
-
-# ----------------------------------------------------------------------------
-# Panel 2 — replication: 32 stored lanes -> 128 reading lanes (4 copies)
-# ----------------------------------------------------------------------------
-ax.text(76, 90, "② Replicated to all 128 warpgroup lanes", ha="center", fontsize=9.6,
- fontweight="bold", color=TXT)
-ax.text(76, 86.3, "R[4 : 32@TLane] — 4 copies at lane stride 32", ha="center", fontsize=7.6,
- color=TXT, style="italic")
-
-# source: the 32 stored lanes
-SX, SY, SW, SH = 52.5, 38, 11, 28
-ax.add_patch(Rectangle((SX, SY), SW, SH, facecolor="#ede9fe", edgecolor="#7c3aed", linewidth=1.8))
-ax.text(SX + SW / 2, SY + SH / 2, "TLane\n0–31\n\n(stored\nonce)", ha="center", va="center",
- fontsize=7.4, fontweight="bold", color="#5b21b6")
-
-# 4 destination quadrants of the reading warpgroup
-DX, DW, DH = 78, 19, 11
-ranges = ["lanes 0–31", "lanes 32–63", "lanes 64–95", "lanes 96–127"]
-dys = [66, 52, 38, 24]
-for i, (rg, dy) in enumerate(zip(ranges, dys)):
- ax.add_patch(Rectangle((DX, dy), DW, DH, facecolor=PURPLE[i], edgecolor="white",
- linewidth=1.4, alpha=0.92))
- ax.text(DX + DW / 2, dy + DH * 0.66, rg, ha="center", va="center", color="white",
- fontsize=7.6, fontweight="bold")
- ax.text(DX + DW / 2, dy + DH * 0.30, "≡ TLane 0–31", ha="center", va="center",
- color="white", fontsize=6.6)
- ax.add_patch(FancyArrowPatch((SX + SW, SY + SH / 2), (DX, dy + DH / 2),
- arrowstyle="-|>", mutation_scale=11,
- color="#8b5cf6", lw=1.3, shrinkA=0, shrinkB=0))
-
-ax.text(76, 16, "lane l reads TLane (l mod 32) — no extra storage.", ha="center",
- fontsize=7.6, color=TXT, style="italic")
-
-# ----------------------------------------------------------------------------
-ax.text(50, 7.5, "Loaded SMEM→TMEM via `tcgen05.cp`; the block-scaled `tcgen05.mma` reads the "
- "scale factors from all 128 warpgroup lanes.", ha="center", fontsize=8.0, color=TXT, style="italic")
-
-fig.savefig(f"{OUT}/sf_tmem.svg", facecolor="white", bbox_inches="tight")
-fig.savefig("/tmp/sf_tmem_preview.png", dpi=130, facecolor="white", bbox_inches="tight")
-plt.close(fig)
-print("wrote sf_tmem.svg")
diff --git a/img/scripts/gen_shuffle_reduce.py b/img/scripts/gen_shuffle_reduce.py
deleted file mode 100644
index 5678858a..00000000
--- a/img/scripts/gen_shuffle_reduce.py
+++ /dev/null
@@ -1,86 +0,0 @@
-"""Generate XOR Shuffle Reduction diagram (chapter_rmsnorm)."""
-import matplotlib
-matplotlib.use('Agg')
-import matplotlib.pyplot as plt
-import matplotlib.patches as mpatches
-import numpy as np
-
-N = 8
-offsets = [4, 2, 1]
-n_steps = len(offsets)
-
-box_w, box_h = 1.2, 0.6
-col_gap = 1.8
-col_stride = box_w + col_gap
-row_stride = 0.85
-
-fig, ax = plt.subplots(figsize=(14, 8))
-ax.axis('off')
-
-thread_colors = ['#4a90d9', '#e67e22', '#27ae60', '#8e44ad',
- '#e74c3c', '#16a085', '#f39c12', '#2c3e50']
-
-values = np.arange(1, N + 1, dtype=float)
-
-def draw_col(x_center, vals, label, is_final=False):
- ax.text(x_center, -0.8, label, ha='center', va='center',
- fontsize=12, fontweight='bold',
- color='#27ae60' if is_final else 'black')
- positions = {}
- for t in range(N):
- y = t * row_stride
- fc = '#d5f5e3' if is_final else '#f0f0f0'
- ec = '#27ae60' if is_final else thread_colors[t]
- ax.text(x_center, y, f"t{t}: {vals[t]:.0f}", ha='center', va='center',
- fontsize=11, fontweight='bold' if is_final else 'normal',
- bbox=dict(boxstyle='round,pad=0.3', fc=fc, ec=ec, lw=2))
- positions[t] = (x_center, y)
- return positions
-
-pos_left = draw_col(0, values, "Initial")
-
-for step_idx, offset in enumerate(offsets):
- new_values = values.copy()
- for t in range(N):
- new_values[t] = values[t] + values[t ^ offset]
-
- is_final = (step_idx == n_steps - 1)
- x_right = (step_idx + 1) * col_stride
- label = "Result" if is_final else f"Step {step_idx}"
- pos_right = draw_col(x_right, new_values, label, is_final)
-
- x_mid = (step_idx * col_stride + x_right) / 2
- ax.text(x_mid, -0.35, f"XOR {offset}", ha='center', va='center',
- fontsize=10, style='italic', color='#666666')
-
- x_src = pos_left[0][0] + box_w / 2
- x_dst = pos_right[0][0] - box_w / 2
-
- for t in range(N):
- partner = t ^ offset
- if partner > t:
- y_t = pos_left[t][1]
- y_p = pos_left[partner][1]
- y_t_r = pos_right[t][1]
- y_p_r = pos_right[partner][1]
- rad = 0.1 + 0.03 * abs(partner - t)
- ax.annotate("", xy=(x_dst, y_t_r), xytext=(x_src, y_p),
- arrowprops=dict(arrowstyle='->', color=thread_colors[partner],
- lw=1.8, connectionstyle=f'arc3,rad={rad}'))
- ax.annotate("", xy=(x_dst, y_p_r), xytext=(x_src, y_t),
- arrowprops=dict(arrowstyle='->', color=thread_colors[t],
- lw=1.8, connectionstyle=f'arc3,rad={rad}'))
-
- values = new_values
- pos_left = pos_right
-
-ax.set_title(f"XOR Shuffle Reduction ({N} threads)\n"
- f"Each step: every thread adds its XOR partner's value. After 3 steps, all threads hold {int(values[0])}.",
- fontsize=13, fontweight='bold', pad=20)
-ax.set_xlim(-1.5, n_steps * col_stride + 1.5)
-ax.set_ylim(-1.3, (N - 1) * row_stride + 0.8)
-ax.invert_yaxis()
-plt.tight_layout()
-plt.savefig('../shuffle_reduce.png', dpi=150, bbox_inches='tight')
-print('Saved shuffle_reduce.png')
-plt.close()
diff --git a/references.bib b/references.bib
deleted file mode 100644
index e69de29b..00000000
diff --git a/setup.py b/setup.py
deleted file mode 100644
index 69ebce97..00000000
--- a/setup.py
+++ /dev/null
@@ -1,28 +0,0 @@
-from setuptools import setup, find_packages
-import tirx_tutorial
-
-requirements = [
- 'jupyter',
- 'numpy',
- 'matplotlib',
- 'requests',
- 'pandas',
- 'sphinx==5.3.0',
- # Pin sphinxcontrib-bibtex and pybtex together: newer sphinxcontrib-bibtex
- # (>=2.6) relies on a private symbol `_FakeEntryPoint` that only exists in
- # pybtex 0.24.x. Mismatched versions break `import sphinxcontrib.bibtex`,
- # which d2lbook always triggers via _build/rst/conf.py.
- 'sphinxcontrib-bibtex<2.6',
- 'pybtex<0.25',
-]
-
-setup(
- name='tirx-tutorial',
- version=tirx_tutorial.__version__,
- python_requires='>=3.10',
- author='MLC Community',
- description='Modern GPU Programming For MLSys',
- packages=find_packages(),
- zip_safe=True,
- install_requires=requirements,
-)
diff --git a/tirx_tutorial/__init__.py b/tirx_tutorial/__init__.py
deleted file mode 100644
index 1064b128..00000000
--- a/tirx_tutorial/__init__.py
+++ /dev/null
@@ -1,5 +0,0 @@
-"""Source code for the TIRx tutorial.
-
-"""
-
-__version__ = "0.0.1"
diff --git a/zh/_extra/demo_zh/barrier_intro.html b/zh/_extra/demo_zh/barrier_intro.html
deleted file mode 100644
index 44a66483..00000000
--- a/zh/_extra/demo_zh/barrier_intro.html
+++ /dev/null
@@ -1,387 +0,0 @@
-
-
-
-
-异步协作：Barrier
-
-
-
-
-
-
-异步协作：Barrier
-硬件单元完成后 signal mbarrier；warp group 在继续执行前等待
-
-
-
-
- GMEM
A, B
-
-
- TMA load
-
- SMEM
-
-
- tcgen05.mma
-
- TMEM
accum
-
-
- tcgen05.ld
-
- Reg
cast
-
-
- store
-
- SMEM
-
-
- TMA store
-
- GMEM
D
-
-
-
-
-
-
-
-
-
-
-
-
-
- TMA
Engine
-
-
- complete-tx
-
-
- Tensor
Core
-
-
-
-
-
-
-
-
-
-
-
-
- - mbarrier (mbar) — 用来协调跨 thread 和硬件单元的异步操作
- - 不同 warp 可以并行执行操作中的不同阶段
-
-
-
-
-
diff --git a/zh/_extra/demo_zh/mbarrier_arrive_timeline.html b/zh/_extra/demo_zh/mbarrier_arrive_timeline.html
deleted file mode 100644
index f3edf659..00000000
--- a/zh/_extra/demo_zh/mbarrier_arrive_timeline.html
+++ /dev/null
@@ -1,470 +0,0 @@
-
-
-
-
-
-MBarrier Arrive 时间线
-
-
-
-
-
- 用于跨 Thread 同步的 MBarrier
-
-
-
-
-
-
-
- T0 (producer)
-
-
-
-
-
-
-
- T1 (consumer)
-
-
-
-
- 时间线
-
-
-
-
-
-
-
-
-
-
-
-
-
- 工作中 / active
-
-
- 发出信号的操作
-
-
- 阻塞 / 等待
-
-
- Barrier 状态
-
-
- 空闲
-
-
- 发信号改变 mbar 状态
-
-
- init 后同步 thread
-
-
-
-
-
- Setup（T0）
- mbarrier.init(mbar, 1) — 所有执行路径只 setup 一次，setup 后需要显式同步。
-
-
- Producer（T0）
- mbarrier.arrive(mbar) 减少 pending_count。当它降到 0 时，phase 完成。
-
-
- Consumer（T1）
- mbarrier.try_wait(mbar, phase) 会挂起 T1，直到 phase 完成。
-
-
-
-
-
-
diff --git a/zh/_extra/demo_zh/mbarrier_tcgen05_timeline.html b/zh/_extra/demo_zh/mbarrier_tcgen05_timeline.html
deleted file mode 100644
index e72a4b22..00000000
--- a/zh/_extra/demo_zh/mbarrier_tcgen05_timeline.html
+++ /dev/null
@@ -1,479 +0,0 @@
-
-
-
-
-
-MBarrier tcgen05.commit 时间线
-
-
-
-
-
- 用 MBarrier 表示 tcgen05 完成
-
-
-
-
-
-
-
- T0 (producer)
-
-
-
- Tensor Core
-
-
-
-
-
-
-
- T1 (consumer)
-
-
-
-
- 时间线
-
-
-
-
-
-
-
-
-
-
-
-
-
- 工作中 / active
-
-
- 发出信号的操作
-
-
- Dispatch 到 Tensor Core
-
-
- 阻塞 / 等待
-
-
- Barrier 状态
-
-
- 空闲
-
-
- 发信号改变 mbar 状态
-
-
- init 后同步 thread
-
-
-
-
-
- Producer（T0）
- tcgen05.mma 发起异步矩阵乘。tcgen05.commit(mbar) 告诉 barrier 跟踪它；完成时硬件会执行 arrive::one。
-
-
- Tensor Core
- 执行 T0 queue 进来的 tcgen05.mma 和 tcgen05.commit。
-
-
- Consumer（T1）
- mbarrier.try_wait(mbar, phase) 会挂起 T1，直到 phase 完成。结果随后可在 TMEM 中读取。
-
-
-
-
-
-