docs: add Mermaid diagrams, Jupyter Notebooks, and course scripts

- docs/diagrams/: 4 Mermaid diagrams (grid-stride loop, thread grid, tree reduction, softmax pipeline)
- notebooks/: 6 Jupyter Notebooks (env setup, copy, vector_add, reduce_sum, softmax, benchmark)
- course/scripts/: 5 course scripts (framework + 4 operators, 5+10+3+2 format)
- fix: JSON encoding in notebooks (BOM removal, Chinese quotation marks)
- fix: troubleshooting.md softmax code example and tilelang-vs-cuda.md template declaration

.github/ISSUE_TEMPLATE/bug_report.md

@@ -0,0 +1,36 @@
+---
+name: Bug Report
+about: 报告代码或文档的错误
+title: "[Bug] "
+labels: bug
+assignees: ""
+---
+
+## 环境信息
+
+- OS: [如 Ubuntu 22.04]
+- CUDA 版本: [如 12.1]
+- GPU 型号: [如 NVIDIA L40]
+- intro-ops 分支/commit: [如 main@abc1234]
+- Python 版本: [如 3.12]
+
+## 描述
+
+请清晰描述你遇到的问题。
+
+## 复现步骤
+
+1. 执行 `...`
+2. 看到错误 `...`
+
+## 期望行为
+
+请描述你期望发生什么。
+
+## 实际行为
+
+请描述实际发生了什么（附上完整的错误信息或截图）。
+
+## 附加信息
+
+如相关：日志、截图、benchmark 数据。

.github/ISSUE_TEMPLATE/feature_request.md

@@ -0,0 +1,23 @@
+---
+name: Feature Request
+about: 提出功能增强或新算子建议
+title: "[Feature] "
+labels: enhancement
+assignees: ""
+---
+
+## 使用场景
+
+请描述你的需求背景：在什么情况下需要这个功能？
+
+## 提议方案
+
+请描述你期望的功能或实现思路。
+
+## 替代方案
+
+是否考虑过其他替代方案？如有请描述。
+
+## 附加信息
+
+如相关：参考实现、论文链接、其他项目的类似功能。

.github/ISSUE_TEMPLATE/question.md

@@ -0,0 +1,22 @@
+---
+name: Question
+about: 使用问题或技术讨论
+title: "[Question] "
+labels: question
+assignees: ""
+---
+
+## 问题描述
+
+请清晰描述你的问题。
+
+## 已尝试的方法
+
+- [ ] 已查阅 README
+- [ ] 已查阅 FAQ
+- [ ] 已搜索 Issues
+- [ ] 已查阅 troubleshooting 指南
+
+## 环境信息（如相关）
+
+- OS / CUDA 版本 / GPU 型号

.github/PULL_REQUEST_TEMPLATE.md

@@ -0,0 +1,26 @@
+## 描述
+
+请简要描述此 PR 做了什么。
+
+## 类型
+
+- [ ] Bug 修复
+- [ ] 新功能
+- [ ] 文档更新
+- [ ] 测试补充
+- [ ] 代码重构
+- [ ] 性能优化
+
+## 验证
+
+- [ ] 相关测试通过：`pytest tests/op_tests/ -v --backend nvidia`
+- [ ] 代码风格检查通过（clang-format / black + isort）
+- [ ] 新增代码有适当的测试覆盖
+
+## 关联 Issue
+
+Closes #
+
+## 截图 / Benchmark（如适用）
+
+（附上相关的测试结果截图或性能数据对比）

course/checklist.md

@@ -0,0 +1,56 @@
+# 学习进度清单
+
+## Phase 1: Kernel 编写
+
+### copy
+- [ ] 理解 grid-stride loop 原理：为什么可以处理任意大小的 tensor
+- [ ] 理解 global memory 合并访问（coalesced access）
+- [ ] NVIDIA `kernel.cuh` TODO 完成
+- [ ] TileLang `kernel.py` TODO 完成
+- [ ] 两种后端测试全部通过
+- [ ] benchmark 跑通，带宽数据合理
+
+### vector_add
+- [ ] 理解逐元素并行的线程网格布局
+- [ ] 理解 tile-level 并行（`T.Parallel` vs `T.Serial` 的区别）
+- [ ] NVIDIA `kernel.cuh` TODO 完成
+- [ ] TileLang `kernel.py` TODO 完成
+- [ ] 两种后端测试全部通过
+- [ ] benchmark 跑通
+
+### reduce_sum（进阶）
+- [ ] 理解 shared memory 树形归约原理
+- [ ] 理解 `__syncthreads()` 的使用时机和条件分支限制
+- [ ] 理解 `T.Serial` 在归约场景中的作用
+- [ ] NVIDIA `kernel.cuh` TODO 完成
+- [ ] TileLang `kernel.py` TODO 完成
+- [ ] 两种后端测试全部通过
+- [ ] 尝试优化 bank conflict
+- [ ] benchmark 跑通
+
+### softmax（挑战）
+- [ ] 理解数值稳定性问题：为什么需要减 max
+- [ ] 理解 online softmax 算法（一遍扫描 vs 三趟扫描）
+- [ ] 理解 log-sum-exp 的滚动更新逻辑
+- [ ] 理解 TileLang 中为什么用 `exp2` / `log2` 替代 `exp` / `log`
+- [ ] NVIDIA `kernel.cuh` TODO 完成
+- [ ] TileLang `kernel.py` TODO 完成
+- [ ] 两种后端测试全部通过
+- [ ] 性能与 PyTorch 参考实现对比
+- [ ] 尝试 warp-level 优化
+
+## Phase 2: 进阶优化
+
+- [ ] 阅读 `docs/tilelang-vs-cuda.md`，理解两种后端的差异
+- [ ] 尝试修改 `.cu` 文件调整 launch 参数（block size、grid size）
+- [ ] 尝试调整 shared memory 大小看性能变化
+- [ ] 学习 CUTLASS/CuTe 集成基础
+- [ ] 尝试提交 benchmark 数据
+
+## Phase 3: 社区贡献
+
+- [ ] 阅读 `docs/how-to-submit-first-pr.md`
+- [ ] 阅读 `docs/how-to-add-an-operator.md`
+- [ ] 找一个 good-first-issue 练手
+- [ ] 提交第一个 PR 并被合并
+- [ ] 尝试贡献文档翻译或 FAQ 补充

course/scripts/00_script_framework.md

@@ -0,0 +1,47 @@
+# 录播/直播课程脚本框架
+
+## 目标
+
+为 intro-ops 四个算子的教学视频提供统一脚本模板。每个视频 **15-20 分钟**，按固定结构组织。
+
+## 统一结构
+
+```
+[5min]  概念讲解：这个算子在深度学习中的用途 + 算法推导
+[10min] 代码实操：IDE 中边写边讲，展示关键决策点
+[3min]  测试验证：跑测试、看 benchmark、分析性能
+[2min]  常见错误演示：故意写错然后排查
+```
+
+## 讲师准备清单
+
+### 课前
+- [ ] 确认录制环境：IDE 字体够大、终端清晰、无弹窗干扰
+- [ ] 打开对应的 kernel 骨架文件（`.cuh` 和 `.py`）
+- [ ] 准备好测试命令（粘贴即用）
+- [ ] 准备一个"故意写错"的版本用于错误演示环节
+
+### 课中
+- [ ] 语速适中（190-210 字/分钟）
+- [ ] 写代码时边写边说——不要沉默
+- [ ] 每个关键决策点停下来解释"为什么"
+- [ ] 错误演示后立即展示正确写法做对比
+
+### 课后
+- [ ] 提供本视频的代码 diff（骨架 → 完整实现）
+- [ ] 在视频描述中附相关文档链接
+- [ ] 设置一个"课后挑战"（如"把带宽从 X 提到 Y"）
+
+## 发布节奏建议
+
+| 平台 | 频率 | 内容类型 |
+|------|------|---------|
+| B 站 | 每周 1 期 | 完整教学视频 |
+| 公众号 | 每周 1 篇 | 视频的文字摘要 + 代码片段 |
+| 知乎 | 每周 1 篇 | 从视频中提炼一个深度问题单独讨论 |
+
+## 字幕与多语言
+
+- 视频字幕优先中文
+- 代码和术语保留英文
+- 关键概念在字幕中标注英文对照（如 "grid-stride loop（网格跨步循环）"）

course/scripts/01_copy_script.md

@@ -0,0 +1,129 @@
+# Copy 算子 — 教学视频脚本
+
+> 时长：15-20 分钟 | 难度：入门
+
+---
+
+## [5min] 概念讲解
+
+### 开场（30s）
+
+"今天我们从 intro-ops 训练营最简单的算子开始——copy。GPU 上做 copy 和 CPU 上不同——你写不出 `memcpy` 那种一行代码，因为 GPU 有成百上千个线程同时在工作。怎么让每个线程知道它该搬哪些数据？这就是 grid-stride loop 要解决的问题。"
+
+### 算子在深度学习中的用途（1min）
+
+"Copy 看起来 trivial，但在深度学习框架中无处不在。Tensor 的 `clone()`、`contiguous()`、数据加载中的 host-to-device 拷贝——底层都是类似的 copy kernel。理解 copy 就是理解 GPU 内存模型的第一步。"
+
+### 算法推导（3.5min）
+
+**关键画面：** 展示 grid-stride-loop Mermaid 图（`docs/diagrams/grid-stride-loop.md`）
+
+讲解要点：
+1. GPU 的线程组织：grid → block → thread
+2. 每个线程如何计算自己的全局索引：`idx = blockIdx.x * blockDim.x + threadIdx.x`
+3. stride = grid 总线程数 = `gridDim.x * blockDim.x`
+4. 为什么需要循环？因为线程数可能少于元素数
+5. 循环条件 `i < N` 保证了越界安全
+
+**过渡语：** "好，理论就这么多。接下来我们在 IDE 里把它写出来。"
+
+---
+
+## [10min] 代码实操
+
+### CUDA Kernel（5min）
+
+**打开文件：** `ops/copy/nvidia/kernel.cuh`
+
+**边写边讲：**
+
+```
+模板参数 T：支持 float、half 等多种数据类型
+__global__：告诉 nvcc 这是 GPU 函数
+命名空间 oprt::copy::nvidia：遵循项目约定
+```
+
+**关键决策点：**
+- "为什么用 `int64_t` 而不是 `int`？——tensor 可能有超过 2^31 个元素"
+- "为什么 `idx` 和 `stride` 的计算不放在循环条件里？——编译器优化和可读性"
+- "为什么循环变量叫 `i` 不是 `idx`？——`idx` 是起始位置，`i` 是当前处理位置"
+
+**写完后展示完整代码。**
+
+### TileLang Kernel（5min）
+
+**打开文件：** `ops/copy/tilelang/kernel.py`
+
+**边写边讲：**
+
+```
+@tilelang.jit：JIT 编译装饰器——和 nvcc 不同，这里编译发生在运行时
+T.Parallel(N // BLOCK_N)：告诉编译器这 N/BLOCK_N 个 tile 可以并行
+T.copy：把全局内存的一整块搬进 fragment
+```
+
+**对比时刻：** "看，TileLang 不需要你写 grid、block、thread——编译器帮你做。这 5 行 Python 生成的 CUDA code 比你手写的可能还高效。"
+
+---
+
+## [3min] 测试验证
+
+### 跑测试
+
+```bash
+# NVIDIA
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
+  pytest tests/op_tests/test_copy.py -v --backend nvidia
+```
+
+**展示：** 终端中测试全部绿色的画面。
+
+### 跑 benchmark
+
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
+  python tests/run_ops.py --op copy --backend nvidia --mode all
+```
+
+**解读 benchmark 输出：**
+- "看这个 Bandwidth——如果你跑不到峰值的 80%+，说明合并访问有问题"
+- "copy 是纯内存搬运，它的性能天花板就是 GPU 内存带宽"
+
+---
+
+## [2min] 常见错误演示
+
+### 错误 1：忘记 grid-stride loop（30s）
+
+```cuda
+// 错误——只处理了前 blockDim.x 个元素
+int idx = threadIdx.x;
+dst[idx] = src[idx];  // 超过 blockDim.x 的部分没被处理！
+```
+
+"这就是为什么需要 grid-stride loop——如果不写循环，每个线程只处理一个元素，你搬不完整个 tensor。"
+
+### 错误 2：循环写反条件（30s）
+
+```cuda
+// 错误——死循环
+for (int i = idx; i < n; i++) {  // 忘记 += stride！
+```
+
+"每次循环 i 都只加 1——第一个线程就把所有活干完了，其他线程白等，而且还死循环。"
+
+### 错误 3：TileLang 里用 T.Serial（1min）
+
+```python
+# 错误——用 T.Serial 跑 copy
+for i in T.Serial(N // BLOCK_N):
+    ...
+```
+
+"T.Serial 告诉编译器'这些 tile 必须顺序执行'——你的并行性全丢了。copy 的每个 tile 完全独立，应该用 T.Parallel。"
+
+---
+
+## 课后挑战
+
+"把 copy kernel 的带宽从 60% 跑到 90% 以上。提示：检查你的 grid/block 尺寸是否合理。下周 vector_add 见！"

course/scripts/02_vector_add_script.md

@@ -0,0 +1,112 @@
+# Vector Add 算子 — 教学视频脚本
+
+> 时长：15-20 分钟 | 难度：入门
+
+---
+
+## [5min] 概念讲解
+
+### 开场（30s）
+
+"上次我们做了 copy——纯内存搬运。今天加一点计算：vector_add，逐元素加法。结构上和 copy 几乎一样，但引入了两个输入和一个辅助函数。"
+
+### 算子在深度学习中的用途（1min）
+
+"逐元素操作（element-wise ops）在深度学习中极常见：ReLU、dropout、residual connection 的加法、batch norm 的 scale+shift。vector_add 是所有这些操作的原型。"
+
+### 算法推导（3.5min）
+
+**关键画面：** 展示线程网格布局图（`docs/diagrams/thread-grid-layout.md`）
+
+讲解要点：
+1. 和 copy 一样的 grid-stride loop 模式
+2. 区别：每次循环处理 `a[i] + b[i]` 而非 `src[i]`
+3. 因为是逐元素，线程间仍然零通信
+4. half 精度需要特殊处理：`__hadd()` 而非 `+`
+
+**过渡语：** "vector_add 的 CUDA kernel 骨架比 copy 多了一个辅助函数——我们来写。"
+
+---
+
+## [10min] 代码实操
+
+### CUDA Kernel（5min）
+
+**打开文件：** `ops/vector_add/nvidia/kernel.cuh`
+
+**Step 1: `add_values<T>` 泛型版（1min）**
+
+```cuda
+template <typename T>
+__device__ T add_values(T a, T b) {
+    return a + b;  // 泛型版本直接用 +
+}
+```
+
+"`__device__` 表示这个函数运行在 GPU 上，只能被 kernel 或其他 device 函数调用。为什么抽出一个函数？因为 half 精度需要特化。"
+
+**Step 2: `add_values<half>` 特化版（1min）**
+
+```cuda
+template <>
+__device__ inline half add_values<half>(half a, half b) {
+    return __hadd(a, b);  // half 专用加法指令
+}
+```
+
+**关键决策点：** "为什么 half 用 `__hadd` 而不是 `+`？CUDA 中 `half` 是存储类型——算术操作需要先转 float 或调用 intrinsics。`__hadd` 是硬件原生指令，更快。"
+
+**Step 3: kernel 函数（3min）**
+
+和 copy kernel 几乎一样，只有行不同：`out[i] = add_values(a[i], b[i]);`
+
+"注意 `__restrict__` 关键字——告诉编译器 `out`、`a`、`b` 指向不重叠的内存区域，可以做更激进的优化。"
+
+### TileLang Kernel（5min）
+
+**打开文件：** `ops/vector_add/tilelang/kernel.py`
+
+**边写边讲：**
+- 外层 `T.Parallel` 分发 tile，内层 `T.Parallel` 分发元素
+- "两层都是 `T.Parallel`——因为所有 (a[i], b[i]) 独立计算"
+
+**对比：** "CUDA 里你手动写了三层抽象（thread → block → grid），TileLang 两层 Parallel 就完成了。编译器把外层的 tile Parallel 映射到 block，内层的 element Parallel 映射到 thread。"
+
+---
+
+## [3min] 测试验证
+
+### 跑测试
+
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
+  pytest tests/op_tests/test_vector_add.py -v --backend nvidia
+```
+
+**重点展示：** "注意测试里包含了 half 精度——如果你 `add_values<half>` 写错了，这里会炸。"
+
+### Benchmark 解读
+
+"vector_add 的理论带宽利用率和 copy 接近——因为它也是 memory-bound 的。如果你的带宽比 copy 低很多，检查是否有非合并访问。"
+
+---
+
+## [2min] 常见错误演示
+
+### 错误 1：half 用 `a + b`（30s）
+
+"CUDA 里 `half + half` 会自动提升为 float 再截断——能编译通过但结果有精度损失。必须用 `__hadd`。"
+
+### 错误 2：不特化 half 就直接在 kernel 里用 `+`（30s）
+
+"`add_values<T>` 里泛型版用 `a + b`，half 特化版用 `__hadd`。如果没写特化版，half 走泛型 `a + b` 也能跑——但性能差。测试能通过，benchmark 会暴露。"
+
+### 错误 3：TileLang 忘记内层 Parallel（1min）
+
+"如果内层用 `T.Serial`——tile 内的 256 个元素就变顺序执行了，线程完全没利用上。结果速度和单线程差不多。"
+
+---
+
+## 课后挑战
+
+"尝试给 vector_add 添加第三个输入：`c = a + b + c`（in-place add）。提示：看 `operator_runtime/ops/vector_add.py` 的 `vector_add_` 函数签名。"

course/scripts/03_reduce_sum_script.md

@@ -0,0 +1,125 @@
+# Reduce Sum 算子 — 教学视频脚本
+
+> 时长：15-20 分钟 | 难度：进阶
+
+---
+
+## [5min] 概念讲解
+
+### 开场（30s）
+
+"前两个算子 copy 和 vector_add，线程之间是零通信的——各自算各自的。今天 reduce_sum 完全不同——线程需要互相通信，通过 shared memory 把部分结果合并为最终结果。"
+
+### 算子在深度学习中的用途（1min）
+
+"归约（reduction）操作在深度学习中非常常见：LayerNorm 里的 mean/std、attention 里的 softmax 分母 sum、loss 函数的最终求和——都是归约。理解归约是 GPU 编程的第一个分水岭。"
+
+### 算法推导（3.5min）
+
+**关键画面：** 展示树形归约流程图（`docs/diagrams/tree-reduction.md`）
+
+讲解要点：
+1. 目标：每行 N 个元素求和为 1 个值
+2. Step 1: 各线程先独立累加自己负责的列 → 存入 shared memory
+3. Step 2: 树形归约——`stride = blockDim.x/2, /4, ..., 1`，每次配对相加
+4. 为什么叫"树形"——每步参与线程减半，log₂(blockDim.x) 步完成
+5. 关键：每一步之后必须 `__syncthreads()`
+
+**可视化辅助：** 手动画出 8→4→2→1 的归约树
+
+**过渡语：** "听起来简单，但写起来有两个坑——我们直接在 IDE 里看。"
+
+---
+
+## [10min] 代码实操
+
+### CUDA Kernel（6min）
+
+**打开文件：** `ops/reduce_sum/nvidia/kernel.cuh`
+
+**Step 1: 线程各自累加（2min）**
+
+"每个线程用 grid-stride loop 跨步累加自己负责的列——这个和 copy 一样。但这次结果不是写到全局内存，而是写到 shared memory。"
+
+**Step 2: shared memory 写入 + 同步（1min）**
+
+```cuda
+smem[threadIdx.x] = sum;
+__syncthreads();  // 关键！
+```
+
+**关键决策点：** "为什么这里一定要 `__syncthreads()`？因为 tree reduction 下一步要读 `smem[tid+s]`——那是别的线程写的。不同步的话可能读到旧数据。"
+
+**Step 3: 树形归约（2min）**
+
+```cuda
+for (int s = blockDim.x / 2; s > 0; s >>= 1) {
+    if (threadIdx.x < s) {
+        smem[threadIdx.x] += smem[threadIdx.x + s];
+    }
+    __syncthreads();  // 在 if 外面！
+}
+```
+
+"注意 `__syncthreads()` 在 if **外面**——这是最常见的 bug。如果放进 if，只有一半线程执行同步，另一半跳过，整个 block 死锁。"
+
+**Step 4: 输出（1min）**
+
+"只有 thread 0 写回全局内存——因为 smem[0] 已经是整行的和。"
+
+### TileLang Kernel（4min）
+
+**打开文件：** `ops/reduce_sum/tilelang/kernel.py`
+
+**边写边讲：**
+- "外层 `T.Parallel` 分发到行——不同行之间独立"
+- "内层 `T.Serial` 顺序遍历列分块——因为累加器有状态依赖"
+- "`T.reduce_sum(tile)` 替代手写树形归约——编译器为你生成最优代码"
+- "注意 `T.alloc_fragment` 分配累加器——这是寄存器级别的存储"
+
+**对比：** "CUDA 版 20+ 行，TileLang 版 8 行。`T.reduce_sum` 内部帮你处理了同步、bank conflict、warp divergence——这些在 CUDA 里都是你要手写的。"
+
+---
+
+## [3min] 测试验证
+
+### 跑测试
+
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
+  pytest tests/op_tests/test_reduce_sum.py -v --backend nvidia
+```
+
+**重点：** "注意测试里的容差——归约涉及很多加法，浮点累加误差比 copy 大。这是正常的，只要在容差内就行。"
+
+### Benchmark 解读
+
+"reduce_sum 的计算量和内存访问量之比（算术强度）比 copy 高——这意味着它从 memory-bound 向 compute-bound 靠近。优化方向也变了：从追求带宽利用率转向减少 bank conflict。"
+
+---
+
+## [2min] 常见错误演示
+
+### 错误 1：`__syncthreads()` 在 if 内（30s）
+
+**故意写错并运行** → 死锁/挂起
+
+"这是最容易踩的坑。症状是程序 hang 住不动。记住：`__syncthreads()` 永远放在条件分支外面。"
+
+### 错误 2：忘记 shared memory 初始化（30s）
+
+"如果把部分和写进 shared memory 之前没有清零——你存的是上一次 kernel launch 的垃圾数据。用 `extern __shared__` 时尤其注意。"
+
+### 错误 3：blockDim.x 不是 2 的幂（30s）
+
+"树形归约假设 block size 是 2 的幂。如果用 300 个线程——stride 从 150 开始，配对就会错位。建议 block size = 128/256/512。"
+
+### 错误 4：写了 shared memory 但忘记声明（30s）
+
+"`extern __shared__ float smem[]` 在 kernel 参数里声明还不够——launch 时 `<<<grid, block, shared_mem_size>>>` 第三个参数必须传。报错 `uses too much shared data` 时检查这里。"
+
+---
+
+## 课后挑战
+
+"消除 reduce_sum 的 bank conflict：给 shared memory 加 padding，对比优化前后的 bandwidth。目标：提升 30%+ 带宽利用率。"

course/scripts/04_softmax_script.md

@@ -0,0 +1,114 @@
+# Softmax 算子 — 教学视频脚本
+
+> 时长：15-20 分钟 | 难度：挑战
+
+---
+
+## [5min] 概念讲解
+
+### 开场（30s）
+
+"softmax 是四个算子中最难的。它结合了 copy 的逐元素、reduce_sum 的归约——还要解决数值溢出问题。今天我们不只写正确的 softmax，还要理解为什么'正确'不是理所当然的。"
+
+### 算子在深度学习中的用途（1min）
+
+"softmax 是 transformer 的核心——每个 attention block 都调它。分类任务的最后一层也是 softmax。一个效率高 10% 的 softmax kernel 可以提升整个推理管线的吞吐。"
+
+### 算法推导（3.5min）
+
+**关键画面：** 展示 softmax 流水线图（`docs/diagrams/softmax-pipeline.md`）
+
+讲解要点：
+
+1. **朴素公式的问题：** `exp(100)` 溢出 → NaN。演示一个小 demo：`torch.tensor([100.0, 200.0, 300.0])` 的朴素 exp 直接炸
+
+2. **稳定的公式：** 
+   - `softmax(x_i) = exp(x_i - max) / Σ exp(x_j - max)`
+   - 为什么等价？分子分母同除 `exp(max)`
+   - 为什么稳定？`x_i - max ≤ 0`，所以 `exp() ≤ 1`，永不溢出
+
+3. **三趟扫描流程：**
+   - Pass 1: 求 `max_val`（reduce max）
+   - Pass 2: 写 `exp(x - max)` 到 out + 累加 sum（reduce sum）
+   - Pass 3: `out[i] /= sum`
+
+4. **Online Softmax（TileLang 版）：**
+   - 用 log-sum-exp (lse) 滚动更新，两趟完成
+   - `exp2`/`log2` 比 `exp`/`log` 在硬件上更快
+
+**过渡语：** "理论比代码复杂——但代码本身并不可怕。来写。"
+
+---
+
+## [10min] 代码实操
+
+### CUDA Kernel（5min）
+
+**打开文件：** `ops/softmax/nvidia/kernel.cuh`
+
+**Step 1: Pass 1 — 求行最大值（1.5min）**
+
+"和 reduce_sum 的归约一模一样——只是把 `+=` 改成 `max()`。注意初始值：max 初始化为 `in[row * cols]` 而不是 0——因为输入可能是全负数。"
+
+**Step 2: Pass 2 — exp + 累加 sum（2min）**
+
+"第二趟扫描做了两件事：计算 `exp(x - max)` 写到 out，同时累加 sum。为什么要写 out？因为第三趟需要这些中间值——不写的话第三趟还得重新从全局内存读输入再算 exp，浪费带宽。"
+
+**Step 3: Pass 3 — 归一化（1.5min）**
+
+"最后一行 `__syncthreads()` 确保 sum 已经归约完成。除了 thread 0 知道 sum 之外——其他线程不需要知道 sum 就能做除法吗？不对——sum 还在 smem[0] 里，每个线程需要读 smem[0] 来做除法。所以第三趟之前也要 sync。"
+
+### TileLang Kernel（5min）
+
+**打开文件：** `ops/softmax/tilelang/kernel.py`
+
+**关键概念：**
+- "`log2_e = 1.44269504`：`exp(x) = 2^(x * log2(e))` 的转换系数"
+- "Pass 1 用 `T.Serial` 遍历列分块——和 reduce_sum 一样，lse 有状态依赖"
+- "`T.reduce_max` + `T.reduce_sum`：编译器为你选择最优的归约策略"
+- "lse 更新公式：`m_new = max(lse, max(tile)); lse = m_new + log2(exp2(lse - m_new) + sum(exp2(tile - m_new)))`"
+
+**不展开完整公式推导——引导学员去看 [docs/diagrams/softmax-pipeline.md](../../docs/diagrams/softmax-pipeline.md) 中的 LSE 滚动更新详解。**
+
+---
+
+## [3min] 测试验证
+
+### 跑测试
+
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia \
+  pytest tests/op_tests/test_softmax.py -v --backend nvidia
+```
+
+**重点展示：** "测试里会验证两件事——正确性（每行和 ≈ 1.0）和数值稳定性（大数据不出 NaN）。如果你偷懒没写减 max，大数值用例会直接炸。"
+
+### Benchmark 解读
+
+"softmax 的算术强度比 reduce_sum 更高——exp/log 是计算密集型操作。优化方向从内存带宽转向计算吞吐。如果你的算力利用率低，可能是 exp 计算没被流水线化。"
+
+---
+
+## [2min] 常见错误演示
+
+### 错误 1：忘记减 max（30s）
+
+"最经典的 bug。输入 `[100, 200, 300]`——朴素 exp 输出全是 NaN。减了 max 后正常输出 `[0, 0, 1]`。"
+
+### 错误 2：Pass 2 和 Pass 3 之间没同步（30s）
+
+"Pass 2 的 sum 归约完成后，thread 0 有正确的 sum——但 thread 1 可能还在写 smem。不 sync 的话 thread 1 在 Pass 3 读到的是旧 sum。"
+
+### 错误 3：max 初始化为 0 而非第一个元素（30s）
+
+"如果输入全是负数——max=0 比真实值大。exp(x - 0) 没问题（仍然 ≤ 1），但 exp(x - real_max) 的精度更好。不影响正确性但影响精度。"
+
+### 错误 4：TileLang 里 Pass 1 用 T.Parallel（30s）
+
+"lse 有跨 tile 的状态依赖——必须 T.Serial。用 T.Parallel 会读到未初始化的 lse 值。"
+
+---
+
+## 课后挑战
+
+"把 CUDA 的三趟扫描改成两趟 online softmax（像 TileLang 版一样）。分析性能提升幅度，写一段注释解释为什么快。"

docs/FAQ.md

@@ -0,0 +1,111 @@
+# 常见问题解答 (FAQ)
+
+> 最后更新：2026-06-05
+
+## 环境配置
+
+### Q: 需要什么硬件才能运行 intro-ops？
+
+**A:** 需要 NVIDIA GPU（支持 CUDA）或沐曦 MetaX GPU。Intel 集显、AMD GPU（无 ROCm 适配）、Apple Silicon 均无法运行。
+
+### Q: 没有 NVIDIA GPU 怎么办？
+
+**A:** 可以考虑租用云端 GPU（AutoDL、恒源云等平台），租一张 T4 或 GTX 1060 即可跑通全部算子。
+
+### Q: CUDA Toolkit 需要什么版本？
+
+**A:** 建议 CUDA 11.8 或以上。CMake 用 `"native"` 架构参数可自动适配你的 GPU。
+
+### Q: 必须用 conda 管理 Python 环境吗？
+
+**A:** 不强制，但推荐。主要依赖是 PyTorch >= 2.0、pytest >= 7.0、cmake >= 3.22。TileLang 后端还需 `pip install tilelang`。
+
+### Q: Windows 上能跑吗？
+
+**A:** 项目设计为 Linux 环境（构建脚本为 bash，产物为 `.so`）。Windows 上建议用 WSL2。
+
+---
+
+## 编译构建
+
+### Q: 构建报 `FindCUDA failed` 怎么办？
+
+**A:** 检查 CUDA Toolkit 是否安装，`nvcc --version` 是否可运行，`CMAKE_CUDA_COMPILER` 是否正确。
+
+### Q: 每次改完 kernel 都要重新构建吗？
+
+**A:** NVIDIA 后端需要重新构建（`bash scripts/build_nvidia.sh build`）。TileLang 后端是 JIT 编译，改完 Python 代码直接跑测试即可。
+
+### Q: `CMake Error: Unknown CUDA architecture` 怎么处理？
+
+**A:** 修改 `CMakePresets.json`，将 `CMAKE_CUDA_ARCHITECTURES` 改为 `"native"` 或你的 GPU 对应架构编号。
+
+### Q: 构建成功但 `import` 时报找不到 `.so` 文件？
+
+**A:** 检查环境变量 `CAMP_BUILD_DIR` 是否指向正确的构建目录（如 `build-nvidia`）。
+
+---
+
+## Kernel 编写
+
+### Q: 四个算子应该按什么顺序学习？
+
+**A:** copy → vector_add → reduce_sum → softmax，难度递增。copy 最简单（纯内存搬运），softmax 需要理解数值稳定性和 online 算法。
+
+### Q: grid-stride loop 为什么能处理任意大小的 tensor？
+
+**A:** 每个线程在循环中处理多个元素（步长为 `gridDim.x * blockDim.x`），而不是只处理一个。这样无论总元素数是多少，只要循环条件 `idx < N` 就能覆盖。
+
+### Q: `__syncthreads()` 什么情况下必须用？
+
+**A:** 当 block 内线程通过 shared memory 交换数据时，写入 shared memory 后必须 `__syncthreads()` 确保所有线程都完成写入，之后才能安全读取。典型场景：reduce_sum 的树形归约每一步之间。
+
+### Q: softmax 为什么要减最大值？
+
+**A:** `exp(88.7) ≈ 1.6e38`，接近 FP32 上限 `3.4e38`。更大的输入值会导致 `exp()` 溢出为 inf。减最大值等价于分子分母同除 `exp(max)`，数学结果不变但数值稳定。
+
+### Q: TileLang 和 CUDA kernel 需要功能完全一致吗？
+
+**A:** 是的，两者应通过相同的测试用例。但实现方式不同——CUDA 控制线程级别，TileLang 控制 tile 级别。
+
+---
+
+## 测试与调试
+
+### Q: 测试报 `tensor not close` 但肉眼看不出来？
+
+**A:** 检查 `rtol`/`atol` 设置。FP16 建议 `rtol=1e-3`，FP32 建议 `rtol=1e-5`。如果使用 TileLang 后端，注意它用 `exp2`/`log2` 而非 `exp`/`log`，会产生细微差异。
+
+### Q: 如何单独跑一个算子的测试？
+
+**A:**
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v --backend nvidia
+```
+
+### Q: 如何同时跑正确性和 benchmark？
+
+**A:**
+```bash
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op copy --backend nvidia --mode all
+```
+
+### Q: 测试结果不稳定怎么办？
+
+**A:** 少量浮点误差波动（尤其是 reduce_sum 这类归约算子）是正常的。如果同一输入每次运行结果差异超过 `rtol=1e-5`（FP32）或 `rtol=1e-3`（FP16），检查是否有未初始化的内存或越界访问。
+
+---
+
+## 贡献流程
+
+### Q: 我可以贡献什么？
+
+**A:** 从 good-first-issue 标签入手比较合适。常见新手任务：修复文档错别字、补充测试用例、翻译文档、添加代码注释。详见 [how-to-submit-first-pr.md](how-to-submit-first-pr.md)。
+
+### Q: PR 需要什么条件才能被合并？
+
+**A:** 需要通过 CI（所有测试通过）、至少一位 reviewer 审核、代码风格符合规范、包含必要的测试。
+
+### Q: 如何添加一个新算子？
+
+**A:** 参考 [how-to-add-an-operator.md](how-to-add-an-operator.md)，完整流程包括 kernel 实现、C API、Python 绑定、测试和 benchmark。

docs/diagrams/grid-stride-loop.md

@@ -0,0 +1,54 @@
+# Grid-Stride Loop 图示
+
+## 概念
+
+Grid-stride loop 是 GPU kernel 处理任意大小 tensor 的核心技术。每个线程不只处理一个元素，而是在循环中负责多个元素（跨步为 grid 总线程数），从而覆盖任意 N。
+
+## Mermaid 图示
+
+```mermaid
+graph TB
+    subgraph "Global Memory (N=10 elements)"
+        M0["[0]"] --- M1["[1]"] --- M2["[2]"] --- M3["[3]"] --- M4["[4]"]
+        M5["[5]"] --- M6["[6]"] --- M7["[7]"] --- M8["[8]"] --- M9["[9]"]
+    end
+
+    subgraph "Grid = 2 Blocks × 2 Threads = 4 Threads"
+        subgraph "Block 0"
+            T0["Thread 0<br/>idx=0"]
+            T1["Thread 1<br/>idx=1"]
+        end
+        subgraph "Block 1"
+            T2["Thread 2<br/>idx=2"]
+            T3["Thread 3<br/>idx=3"]
+        end
+    end
+
+    T0 -->|"i=0"| M0
+    T0 -->|"i=4"| M4
+    T0 -->|"i=8"| M8
+    T1 -->|"i=1"| M1
+    T1 -->|"i=5"| M5
+    T1 -->|"i=9"| M9
+    T2 -->|"i=2"| M2
+    T2 -->|"i=6"| M6
+    T3 -->|"i=3"| M3
+    T3 -->|"i=7"| M7
+```
+
+## 代码对应
+
+```cuda
+// stride = gridDim.x * blockDim.x = 2 * 2 = 4
+for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x;  // idx = 0,1,2,3
+     i < N;                                               // N = 10
+     i += gridDim.x * blockDim.x) {                       // i += 4
+    dst[i] = src[i];
+}
+```
+
+## 要点
+
+- 每个线程负责多个元素，间隔为 stride（本例中 stride=4）
+- 无论 N 多大，只要 `i < N` 条件在，就不会越界
+- grid/block 尺寸可自由调整，不影响正确性，只影响性能

docs/diagrams/softmax-pipeline.md

@@ -0,0 +1,66 @@
+# Softmax 流水线对比
+
+## 朴素三趟扫描 vs Online Softmax
+
+### 朴素实现（三趟扫描）
+
+```
+Pass 1: 求 row_max        → max_val
+Pass 2: 求 exp(x-max) 和  → sum (写入中间结果到 out)
+Pass 3: 除以 sum           → 最终归一化结果
+```
+
+### Online Softmax（两趟扫描）
+
+```
+Pass 1: 滚动更新 lse (log-sum-exp) → 不写 out
+Pass 2: 用最终 lse 归一化          → 写入最终结果
+```
+
+## 流程图对比
+
+```mermaid
+flowchart LR
+    subgraph "朴素三趟扫描"
+        direction TB
+        N1["输入 x[0..N-1]"]
+        N2["Pass 1: 遍历全行<br/>max = max(x)"]
+        N3["Pass 2: 遍历全行<br/>sum = Σ exp(x - max)<br/>out[i] = exp(x[i]-max)"]
+        N4["Pass 3: 遍历全行<br/>out[i] /= sum"]
+        N5["输出 out[0..N-1]"]
+        N1 --> N2 --> N3 --> N4 --> N5
+    end
+
+    subgraph "Online Softmax（两趟扫描）"
+        direction TB
+        O1["输入 x[0..N-1]"]
+        O2["Pass 1: 滚动 lse<br/>lse = -∞<br/>for chunk in x:<br/>  m_new = max(lse, max(chunk))<br/>  lse = m_new + log(exp(lse-m_new)·out_old + Σ exp(chunk-m_new))"]
+        O3["Pass 2: 归一化<br/>for chunk in x:<br/>  out[i] = exp2(log2_e · (x[i] - lse))"]
+        O4["输出 out[0..N-1]"]
+        O1 --> O2 --> O3 --> O4
+    end
+```
+
+## LSE 滚动更新详解
+
+```mermaid
+flowchart TD
+    Init["lse = -∞"]
+    Chunk1["处理 chunk 1<br/>m_new = max(-∞, max(chunk1)) = max(chunk1)<br/>lse = m_new + log(0 + Σ exp(chunk1 - m_new))"]
+    Chunk2["处理 chunk 2<br/>m_new = max(lse, max(chunk2))<br/>如果 m_new > lse: 旧结果需要 rescale<br/>lse = m_new + log(exp(lse-m_new) + Σ exp(chunk2-m_new))"]
+    Final["最终 lse 包含全行的 log-sum-exp"]
+
+    Init --> Chunk1 --> Chunk2 --> Final
+```
+
+## 要点
+
+| 方面 | 朴素 | Online |
+|------|------|--------|
+| 扫描次数 | 3 趟 | 2 趟 |
+| 全局内存写 | 2 次 (中间结果 + 最终) | 1 次 (仅最终) |
+| 数值稳定性 | 靠减 max 保证 | 靠滚动 lse 保证 |
+| TileLang 中的实现 | N/A | 用 `exp2`/`log2` 替代 `exp`/`log` |
+
+- TileLang 版本使用 `exp2`/`log2` 因为硬件对 2 的幂运算支持更好
+- `log2_e = 1.44269504` 是转换系数：`exp(x) = 2^(x * log2(e))`

docs/diagrams/thread-grid-layout.md

@@ -0,0 +1,71 @@
+# 线程网格布局 — vector_add
+
+## 概念
+
+vector_add 是最经典的逐元素并行模式。输入 a 和 b 是两个等长一维 tensor，输出 c[i] = a[i] + b[i]。每个线程处理一组独立元素，线程之间无需通信。
+
+## Mermaid 图示
+
+```mermaid
+graph TB
+    subgraph "输入 A (N=8)"
+        A0["a[0]"] --- A1["a[1]"] --- A2["a[2]"] --- A3["a[3]"]
+        A4["a[4]"] --- A5["a[5]"] --- A6["a[6]"] --- A7["a[7]"]
+    end
+
+    subgraph "输入 B (N=8)"
+        B0["b[0]"] --- B1["b[1]"] --- B2["b[2]"] --- B3["b[3]"]
+        B4["b[4]"] --- B5["b[5]"] --- B6["b[6]"] --- B7["b[7]"]
+    end
+
+    subgraph "Grid (4 Threads, stride=4)"
+        T0["Thread 0<br/>i=0,4"]
+        T1["Thread 1<br/>i=1,5"]
+        T2["Thread 2<br/>i=2,6"]
+        T3["Thread 3<br/>i=3,7"]
+    end
+
+    subgraph "输出 C (N=8)"
+        C0["c[0]"] --- C1["c[1]"] --- C2["c[2]"] --- C3["c[3]"]
+        C4["c[4]"] --- C5["c[5]"] --- C6["c[6]"] --- C7["c[7]"]
+    end
+
+    A0 --> T0 --> C0
+    B0 --> T0
+    A4 --> T0 --> C4
+    B4 --> T0
+    A1 --> T1 --> C1
+    B1 --> T1
+    A5 --> T1 --> C5
+    B5 --> T1
+    A2 --> T2 --> C2
+    B2 --> T2
+    A6 --> T2 --> C6
+    B6 --> T2
+    A3 --> T3 --> C3
+    B3 --> T3
+    A7 --> T3 --> C7
+    B7 --> T3
+```
+
+## 代码对应
+
+```cuda
+template <typename T>
+__global__ void vector_add_contiguous_kernel(
+    T *out, const T *a, const T *b, int64_t n) {
+
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    int64_t stride = gridDim.x * blockDim.x;
+
+    for (int64_t i = idx; i < n; i += stride) {
+        out[i] = a[i] + b[i];
+    }
+}
+```
+
+## 要点
+
+- 每个 `(a[i], b[i])` 对独立计算，线程间零通信
+- 相邻线程访问相邻内存地址 → 合并访问（coalesced access）
+- TileLang 版本用 `T.Parallel` 表达同样的并行语义

docs/diagrams/tree-reduction.md

@@ -0,0 +1,85 @@
+# 树形归约 — reduce_sum
+
+## 概念
+
+reduce_sum 对矩阵每行做求和归约。一个 block 处理一行，block 内线程先将各自负责的列累加为部分和，再通过 shared memory 做树形归约，最终 thread 0 输出整行的和。
+
+## 树形归约过程（Block Size = 8）
+
+```mermaid
+flowchart TD
+    subgraph "Step 0: 初始化 Shared Memory"
+        S0["smem[0]=3"] --- S1["smem[1]=1"] --- S2["smem[2]=7"] --- S3["smem[3]=0"]
+        S4["smem[4]=4"] --- S5["smem[5]=1"] --- S6["smem[6]=6"] --- S7["smem[7]=3"]
+    end
+
+    subgraph "Step 1: stride=4"
+        A0["smem[0] += smem[4] → 7"] -..- A1["smem[1] += smem[5] → 2"]
+        A2["smem[2] += smem[6] → 13"] -..- A3["smem[3] += smem[7] → 3"]
+    end
+
+    subgraph "Step 2: stride=2"
+        B0["smem[0] += smem[2] → 20"] -..- B1["smem[1] += smem[3] → 5"]
+    end
+
+    subgraph "Step 3: stride=1"
+        C0["smem[0] += smem[1] → 25"]
+    end
+
+    S0 --> A0
+    S4 --> A0
+    S1 --> A1
+    S5 --> A1
+    S2 --> A2
+    S6 --> A2
+    S3 --> A3
+    S7 --> A3
+    A0 --> B0
+    A2 --> B0
+    A1 --> B1
+    A3 --> B1
+    B0 --> C0
+    B1 --> C0
+
+    C0 --> Result["out[row] = smem[0] = 25"]
+```
+
+## 代码对应
+
+```cuda
+__global__ void reduce_sum_rowwise_kernel(
+    float *out, const float *in, int64_t rows, int64_t cols) {
+
+    extern __shared__ float smem[];
+    int row = blockIdx.x;
+    int tid = threadIdx.x;
+
+    // 1. 各线程累加自己负责的列
+    float sum = 0.0f;
+    for (int c = tid; c < cols; c += blockDim.x) {
+        sum += in[row * cols + c];
+    }
+    smem[tid] = sum;
+    __syncthreads();
+
+    // 2. 树形归约
+    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            smem[tid] += smem[tid + s];
+        }
+        __syncthreads();
+    }
+
+    // 3. 输出
+    if (tid == 0) {
+        out[row] = smem[0];
+    }
+}
+```
+
+## 要点
+
+- 每步 stride 减半，log₂(blockDim.x) 步完成
+- 每步之间必须 `__syncthreads()` —— 确保所有线程的 shared memory 写入对下一步可见
+- `__syncthreads()` **不能**放在 `if (tid < s)` 分支内，否则死锁
+- 优化方向：bank conflict 消除（添加 padding）

docs/en/phase1-kernel-writing.md

@@ -0,0 +1,106 @@
+# Phase 1 Training Objective: Writing Kernels
+
+## Goal
+
+This phase has one primary task:
+
+**Write kernel files for the existing operators.**
+
+All other code (descriptor lifecycle, Python bindings, tests, build system) is already implemented. Your only job is to fill in the computation logic.
+
+Default requirements:
+
+1. **Priority: get the kernel correct and running.**
+2. You do NOT need to pursue peak performance initially.
+3. Complete the TODOs in `kernel.cuh` / `kernel.py` first, then consider aggressive optimizations.
+
+If you've completed the basic goals and want to push performance further, you can modify the wrapper layer (e.g., `.cu` files for the NVIDIA backend), adjust launch policy, switch between kernels, or add more aggressive specializations. These are advanced topics and not required for Phase 1.
+
+---
+
+## Files to Write
+
+Each operator has two backends, corresponding to two kernel files:
+
+| Operator | NVIDIA kernel | TileLang kernel |
+|----------|--------------|-----------------|
+| `copy` | `ops/copy/nvidia/kernel.cuh` | `ops/copy/tilelang/kernel.py` |
+| `vector_add` | `ops/vector_add/nvidia/kernel.cuh` | `ops/vector_add/tilelang/kernel.py` |
+| `reduce_sum` | `ops/reduce_sum/nvidia/kernel.cuh` | `ops/reduce_sum/tilelang/kernel.py` |
+| `softmax` | `ops/softmax/nvidia/kernel.cuh` | `ops/softmax/tilelang/kernel.py` |
+
+Recommended order (easy to hard): copy → vector_add → reduce_sum → softmax.
+
+---
+
+## NVIDIA Kernel (`.cuh` files)
+
+### What to Write
+
+A `__global__` function placed under the operator's namespace. The function handles only computation logic, taking raw pointers as input/output. It does NOT involve any descriptor or public API.
+
+By default, you do NOT need to modify the `.cu` files in the same directory this phase. The repository already provides compatibility-focused launch boilerplate. This phase only requires you to make `kernel.cuh` correct and runnable.
+
+If you've completed the basic goal and want to push performance, treat the `.cu` file as an advanced optimization layer: adjust thread counts there, switch kernel variants, or add specialized paths. These are outside the scope of Phase 1.
+
+### Difficulty Progression
+
+**copy / vector_add**: Element-wise operations using the grid-stride loop pattern. Each thread handles several independent elements with no inter-thread communication.
+
+**reduce_sum**: Row-wise reduction. One block per row. Threads within the block first independently accumulate their assigned columns, then perform a tree reduction via shared memory. `__syncthreads()` is required for synchronization.
+
+**softmax**: The basic version uses a three-pass flow: first pass finds the row max (numerical stability), second pass computes `exp(x - max)` sum and writes intermediate results to output, third pass divides by the sum. Each pass requires intra-block synchronization.
+
+### Concepts to Understand
+
+- **grid-stride loop**: Why this pattern handles arbitrarily-sized tensors
+- **shared memory reduction**: What each step of tree reduction does, and why `__syncthreads()` is necessary
+- **softmax minus max**: Why computing `exp(x)` directly causes problems, and why subtracting the row max does not change the result
+
+---
+
+## TileLang Kernel (`.py` files)
+
+### What to Write
+
+A Python function decorated with `@tilelang.jit` that describes tile-level computation using the TileLang DSL. TileLang compiles it into real CUDA kernels.
+
+As with NVIDIA, you only need to complete the TODOs in `kernel.py` this phase — no need to modify outer adapter code.
+
+### Difficulty Progression
+
+**copy**: Use the built-in `T.copy` to move data between two tiles. No manual loop needed.
+
+**vector_add**: Use a `T.Parallel` loop to compute element-wise within a tile. Understand the difference between `T.Parallel` and `T.Serial`.
+
+**reduce_sum**: Block-wise accumulation required. The outer loop uses `T.Serial` to traverse column-direction blocks sequentially (because state accumulates), and the inner loop uses `T.reduce_sum` to reduce over a fragment.
+
+**softmax**: Two-pass scan using the online softmax algorithm. The first pass maintains a rolling log-sum-exp state; the second pass normalizes with the final LSE. Uses `exp2` / `log2` instead of `exp` / `log`.
+
+### Concepts to Understand
+
+- **`T.Parallel` vs `T.Serial`**: When can loop iterations run in parallel, and when must they be sequential
+- **`T.alloc_fragment`**: A tile-level local buffer, corresponding to registers or shared memory
+- **`T.copy`**: Moves a chunk of global memory into a fragment — not element-by-element assignment
+- **online softmax**: Why a single pass can produce the correct normalization, and the rolling update logic for log-sum-exp
+- **`exp2` / `log2`**: Why TileLang uses these instead of `exp` / `log`
+
+---
+
+## Verification
+
+After writing each kernel, verify with the corresponding test:
+
+```bash
+# NVIDIA kernel (requires rebuild)
+bash scripts/build_nvidia.sh build
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v --backend nvidia
+
+# TileLang kernel (no build needed)
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_<op>.py -v --backend tilelang
+
+# Both correctness + benchmark
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op <op> --backend nvidia --mode all
+```
+
+All four operators passing both backends = Phase 1 complete.

docs/glossary.md

@@ -0,0 +1,60 @@
+# 术语表 (Glossary)
+
+## GPU 编程
+
+| 英文 | 中文 | 说明 |
+|------|------|------|
+| operator / op | 算子 | GPU 上执行的计算单元，如 copy、softmax |
+| kernel | 核函数 | 在 GPU 上并行执行的函数，由 host 端调用 |
+| grid | 网格 | CUDA 中最高层级的线程组织单元，由多个 block 组成 |
+| block | 线程块 | grid 的子单元，block 内线程可以通过 shared memory 通信 |
+| thread | 线程 | GPU 上最小的执行单元 |
+| warp | 线程束 | GPU 调度的最小执行单元（NVIDIA 上为 32 个线程） |
+| grid-stride loop | 网格跨步循环 | 让 kernel 处理任意大小输入的技术：每个线程负责多个元素，跨度为 grid 总线程数 |
+| shared memory | 共享内存 | block 内线程共享的片上内存（SMEM），速度远超全局内存 |
+| global memory | 全局内存 | GPU 的片外显存（HBM/DRAM），容量大但延迟高 |
+| register | 寄存器 | 每个线程私有的最快存储 |
+| SM (Streaming Multiprocessor) | 流式多处理器 | GPU 上的计算核心单元，一个 GPU 包含多个 SM，每个 SM 可并行执行多个 warp |
+| coalesced access | 合并访问 | 相邻线程访问相邻地址，GPU 可将多次访问合并为一次内存事务 |
+| stride | 步长 | tensor 在某个维度上相邻元素之间的内存地址间隔（以元素数计） |
+| contiguous | 连续 | 描述 tensor 在内存中是否紧密排列、无间隙，合并访问的前提条件 |
+| bank conflict | 存储体冲突 | shared memory 中多个线程同时访问同一 bank 的不同地址导致串行化 |
+| __syncthreads() | 线程同步 | block 内所有线程到达此处后才能继续，用于 shared memory 写入后同步 |
+| reduction | 归约 | 将多个值合并为一个值的操作（求和、求最大等） |
+| online softmax | 在线 softmax | 一遍扫描完成 softmax 的算法，无需单独求 max 和 sum |
+| occupancy | 占用率 | 每个 SM 上活跃 warp 数与理论最大值的比值，影响隐藏延迟的能力 |
+| nvcc | NVIDIA CUDA 编译器 | 将 .cu 文件编译为 GPU 可执行的二进制 |
+| CUDA | — | NVIDIA 的 GPU 通用并行计算平台与编程模型 |
+
+## 数值精度
+
+| 英文 | 中文 | 说明 |
+|------|------|------|
+| FP32 / float32 | 单精度浮点 | 32 位浮点数，约 7 位有效数字 |
+| FP16 / float16 / half | 半精度浮点 | 16 位浮点数，约 3 位有效数字，显存占用为 FP32 的一半 |
+| inf / NaN | 无穷/非数 | 浮点溢出或无效运算的结果 |
+| tolerance / rtol / atol | 容差 | 测试中允许的数值误差上限 |
+
+## 框架术语
+
+| 英文 | 中文 | 说明 |
+|------|------|------|
+| descriptor | 描述符 | 封装算子运行所需全部信息（输入输出 shape、dtype、参数）的对象 |
+| tensor_view | 张量视图 | 对 tensor 的轻量级引用，包含数据指针、shape、stride、dtype |
+| backend | 后端 | 算子的具体硬件实现（NVIDIA CUDA、MetaX MACA、TileLang） |
+| workspace | 工作空间 | kernel 执行时需要的临时内存 |
+| launch / dispatch | 启动/分发 | 将 kernel 函数提交到 GPU 执行队列 |
+| FFI | 外部函数接口 | Python 调用 C 动态库的机制（本项目用 ctypes） |
+
+## TileLang 术语
+
+| 英文 | 中文 | 说明 |
+|------|------|------|
+| TileLang | — | Python DSL，用 tile 级别的语义描述计算，JIT 编译为 GPU kernel |
+| fragment | 片段 | tile 在单个线程上的局部数据块，对应寄存器或 shared memory |
+| T.alloc_fragment | 分配片段 | 创建 tile 级别的局部 buffer |
+| T.copy | 拷贝 | 在 tile 和全局内存之间搬移数据 |
+| T.Parallel | 并行循环 | 循环的每次迭代可并行执行（不依赖前一次迭代的结果） |
+| T.Serial | 顺序循环 | 循环必须按序执行（后一次依赖前一次的结果） |
+| T.reduce_sum | 归约求和 | 对 fragment 内的元素做求和规约 |
+| JIT | 即时编译 | 运行时将 TileLang Python 函数编译为 GPU 机器码 |

docs/how-to-submit-first-pr.md

@@ -0,0 +1,154 @@
+# 如何为 intro-ops 提交第一个 PR
+
+## 第一步：环境准备
+
+### Fork 仓库 + Clone
+
+1. 打开 https://github.com/metax/intro-ops ，点击右上角 **Fork** 按钮
+2. Clone 你 fork 的仓库：
+   ```bash
+   git clone https://github.com/<你的用户名>/intro-ops.git
+   cd intro-ops
+   ```
+3. 添加上游仓库（保持同步）：
+   ```bash
+   git remote add upstream https://github.com/metax/intro-ops.git
+   ```
+
+### 配置开发环境
+
+```bash
+# 创建 conda 环境
+conda create -n intro-ops python=3.12 -y
+conda activate intro-ops
+
+# 安装 Python 依赖
+pip install -r requirements.txt
+```
+
+### 构建并验证基线
+
+```bash
+# NVIDIA 后端构建
+bash scripts/build_nvidia.sh configure
+bash scripts/build_nvidia.sh build
+
+# 跑一遍全部测试，确保基线通过
+PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia python tests/run_ops.py --op all --backend nvidia --mode test
+```
+
+如果基线测试全部通过，说明环境配置正确，可以开始贡献。
+
+---
+
+## 第二步：选择你的贡献
+
+在仓库的 [Issues](https://github.com/metax/intro-ops/issues) 页面，按标签筛选适合你的任务：
+
+### 推荐的新手任务（good-first-issue）
+
+| 类型 | 预计耗时 | 示例 |
+|------|---------|------|
+| 修复文档错别字/格式问题 | 15-30 min | 修正 README 中的拼写错误 |
+| 补充测试用例 | 1-2 h | 为某个算子增加 corner case 测试 |
+| 添加代码注释 | 30 min | 给关键算法步骤补充注释 |
+| 翻译文档段落 | 1-2 h | 将中文文档翻译为英文（或反过来） |
+| 补充 FAQ 条目 | 30 min | 把你遇到并解决的问题写成 FAQ |
+
+### 不知道怎么选？
+
+- 在 Issue 下留言询问
+- 或者先读一遍文档，过程中发现任何可改进的地方记录下来
+- 首次贡献从最小的事情做起，熟悉流程比贡献大小更重要
+
+---
+
+## 第三步：开发和自测
+
+### 创建分支
+
+```bash
+git checkout -b fix/your-description
+# 分支命名建议：
+# fix/xxx   — 修复
+# feat/xxx  — 新功能
+# docs/xxx  — 文档
+# test/xxx  — 测试
+```
+
+### 修改代码并本地验证
+
+- **文档类**：修改后确认 Markdown 渲染正常
+- **代码类**：修改后跑相关测试确保不引入回归：
+  ```bash
+  PYTHONPATH=python:. CAMP_BUILD_DIR=build-nvidia pytest tests/op_tests/test_copy.py -v --backend nvidia
+  ```
+
+### 代码风格
+
+- C++ 代码：用 `clang-format` 格式化
+- Python 代码：用 `black` + `isort` 格式化
+
+---
+
+## 第四步：提交 PR
+
+### Commit 规范
+
+```bash
+git add <修改的文件>
+git commit -m "docs: fix typo in README"
+# commit message 格式：<类型>: <简短描述>
+# 类型：docs / fix / feat / test / refactor
+```
+
+### 推送并创建 PR
+
+```bash
+git push origin fix/your-description
+```
+
+然后在 GitHub 上打开你的仓库，点击 **Compare & pull request** 按钮。
+
+### 填写 PR 描述
+
+一个好的 PR 描述应包含：
+1. **做了什么**（一句话概括）
+2. **为什么做**（问题背景）
+3. **如何验证**（跑过哪些测试、附上截图或测试输出）
+4. **关联 Issue**（如有，用 `Closes #123` 关联）
+
+---
+
+## 第五步：响应 Review
+
+1. Review 意见是改进代码的机会，不要抵触
+2. 如有不理解处，在 PR 评论区直接提问
+3. 修改后 push 到同一分支，PR 会自动更新
+4. 所有 Review 意见解决后，reviewer 会合并
+
+---
+
+## 合并之后
+
+- 你的名字将出现在仓库的贡献者列表中
+- 可以把你解决问题的经历写成经验分享（blog 或 Discussions）
+- 继续探索更多贡献类型，挑战更难的任务
+
+---
+
+## 常见问题
+
+**Q: 我的 PR 迟迟没人 review 怎么办？**
+A: 通常 48h 内会有人响应。如果超时，可以在 PR 评论区 @ 仓库维护者。
+
+**Q: 我不确定我的修改是否有价值？**
+A: 可以先开一个 Issue 描述你的想法，得到确认后再动手。
+
+**Q: 合并后有冲突怎么办？**
+A: 将上游 main 合并到你的分支：
+```bash
+git fetch upstream
+git merge upstream/main
+# 解决冲突后 git push
+```

docs/learning-path.md

@@ -0,0 +1,121 @@
+# 分层学习路径
+
+intro-ops 的四个算子从 copy 到 softmax 难度渐进。本文提供三级学习路线，帮你根据自己的基础和时间选择合适的路径。
+
+---
+
+## 入门级（约 1-2 周）
+
+**目标：** 跑通 copy + vector_add，建立 GPU 编程基本概念
+
+**前置知识：**
+- C++ 基础（指针、模板、循环）
+- 基本 GPU 概念：什么是 grid / block / thread
+
+**推荐资源：**
+- 《CUDA C Programming Guide》第 1-3 章（NVIDIA 官方，免费）
+- intro-ops 的 `docs/glossary.md` 术语表
+
+**学习路径：**
+
+1. 阅读 `docs/phase1-kernel-writing.md` 了解整体目标
+2. 阅读 `docs/glossary.md` 掌握基础术语
+3. **copy 算子**：
+   - 理解 grid-stride loop（为什么这样写能处理任意大小 tensor）
+   - 完成 `ops/copy/nvidia/kernel.cuh` 的 TODO
+   - 完成 `ops/copy/tilelang/kernel.py` 的 TODO
+   - 跑通测试：`pytest tests/op_tests/test_copy.py -v --backend nvidia`
+4. **vector_add 算子**：
+   - 理解逐元素并行和 `T.Parallel`
+   - 完成 `ops/vector_add/nvidia/kernel.cuh` 的 TODO
+   - 完成 `ops/vector_add/tilelang/kernel.py` 的 TODO
+   - 跑通测试
+
+**产出：** 两个算子两种后端全部通过测试
+
+**如果你卡住了：** 查看 `docs/troubleshooting.md` 的"运行时错误"章节
+
+---
+
+## 进阶级（约 2-4 周）
+
+**目标：** 完成 reduce_sum + softmax，理解 shared memory 编程
+
+**前置知识：**
+- 已完成入门级
+- 了解 shared memory 概念
+- 了解线程同步机制
+
+**推荐资源：**
+- 《Professional CUDA C Programming》shared memory 章节
+- `docs/tilelang-vs-cuda.md` 对比教程
+
+**学习路径：**
+
+1. **reduce_sum 算子**：
+   - 理解 shared memory 树形归约原理
+   - 理解 `__syncthreads()` 的使用时机（为什么不能放在条件分支内）
+   - 完成 `ops/reduce_sum/nvidia/kernel.cuh` 的 TODO
+   - 完成 `ops/reduce_sum/tilelang/kernel.py` 的 TODO（理解 `T.Serial` 的作用）
+   - 跑通测试
+   - 尝试优化 bank conflict（加 padding）
+
+2. **softmax 算子**：
+   - 理解数值稳定性：为什么要减 max
+   - 理解 online softmax 算法（一遍扫描 vs 三趟扫描）
+   - 完成 `ops/softmax/nvidia/kernel.cuh` 的 TODO
+   - 完成 `ops/softmax/tilelang/kernel.py` 的 TODO（理解 `exp2`/`log2` 的用法）
+   - 跑通测试
+
+3. 提交 benchmark 数据，对比自己的实现与 PyTorch 的差异
+
+**产出：** 全部四个算子两种后端通过测试 + benchmark 数据
+
+---
+
+## 专精级（约 5 周+）
+
+**目标：** 性能调优 + 贡献新算子，从"会用"到"能写"
+
+**前置知识：**
+- 已完成进阶级
+- 对 GPU 架构有一定了解（SM、warp、memory hierarchy）
+
+**推荐资源：**
+- CUTLASS 官方文档
+- `docs/how-to-add-an-operator.md` 新增算子指南
+
+**可选方向：**
+
+### 方向 A：性能调优
+- 分析 benchmark 数据，找到瓶颈（内存带宽 / 计算 / 延迟）
+- 调整 launch 参数（block size、grid size）
+- 尝试 warp-level shuffle 优化
+- 将你的优化写成经验分享
+
+### 方向 B：CUTLASS / CuTe 集成
+- 将 CUTLASS 的 tiling / warp-level 优化应用到现有算子
+- 理解 CuTe 的 layout 和 tensor 抽象
+
+### 方向 C：贡献新算子
+- 阅读 `docs/how-to-add-an-operator.md`
+- 选一个感兴趣的算子（如 LayerNorm、GELU、attention）
+- 按流程实现：kernel → C API → Python 绑定 → 测试 → benchmark
+- 提交 PR
+
+### 方向 D：社区贡献
+- 阅读 `docs/how-to-submit-first-pr.md`
+- 翻译文档（中→英或英→中）
+- 补充 FAQ / troubleshooting 条目
+- Review 其他人的 PR
+
+**产出：** PR 被合并 / 新算子通过 review / 在社区分享经验
+
+---
+
+## 学习建议
+
+- **不要一次追求完美**：先写对再优化，能跑通的 kernel 比跑不快的 kernel 好一万倍
+- **善用 checklist**：`course/checklist.md` 帮你追踪进度
+- **遇到问题先查 FAQ 和 troubleshooting**：大部分常见问题已经有答案
+- **把你的经验写下来**：帮你解决问题的经历，也可能是别人的 FAQ 条目

docs/tilelang-vs-cuda.md

@@ -0,0 +1,163 @@
+# TileLang vs CUDA 对比教程
+
+intro-ops 的特色之一是同时支持 CUDA 和 TileLang 两种后端。本文通过 copy 算子的两种实现，帮助理解"同一个算法，两种 DSL 的表达差异"。
+
+## 概述对比
+
+| 维度 | CUDA | TileLang |
+|------|------|----------|
+| 抽象层级 | 线程级（thread/warp/block） | tile 级（fragment/tile） |
+| 编写难度 | 需关注同步、shared memory 细节 | 声明式，编译器处理调度 |
+| 编译方式 | nvcc 编译 `.cu` 文件 | JIT 编译，无需构建 |
+| 调优手段 | 手动调整 grid/block/shared memory | `T.Parallel` / `T.Serial` 语义 |
+| 适用场景 | 极致性能调优 | 快速原型 + 跨硬件 |
+| 学习曲线 | 陡峭 | 平缓 |
+
+---
+
+## 案例一：copy 算子
+
+### CUDA 实现
+
+```cuda
+template <typename T>
+__global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {
+    // 1. 计算全局线程索引
+    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
+    // 2. 计算网格总步长
+    int64_t stride = gridDim.x * blockDim.x;
+    // 3. grid-stride loop
+    for (int64_t i = idx; i < n; i += stride) {
+        dst[i] = src[i];
+    }
+}
+```
+
+**关键细节：**
+- 你需要手动管理线程索引（`blockIdx.x`、`blockDim.x`、`threadIdx.x`）
+- 你需要手动设计 grid-stride loop 来处理任意大小的 tensor
+- launch 时你需要自己指定 grid size 和 block size
+- 你需要理解"一个线程负责哪些元素"的映射关系
+
+### TileLang 实现
+
+```python
+@tilelang.jit
+def copy_kernel(src, BLOCK_N: int, dtype):
+    N = T.const("N")
+    src: T.Tensor((N,), dtype)
+    out = T.empty((N,), dtype)
+
+    for i in T.Parallel(N // BLOCK_N):
+        tile = T.copy(src[i * BLOCK_N : (i + 1) * BLOCK_N])
+        out[i * BLOCK_N : (i + 1) * BLOCK_N] = tile
+
+    return out
+```
+
+**关键细节：**
+- `T.Parallel` 声明一个 tile 级别的并行循环——编译器负责把 tile 分配到线程
+- `T.copy` 加载一个 tile 到局部 fragment——不需要手写线程索引
+- 没有显式的 grid/block/thread 概念
+- 编译器自动决定最优的 launch 参数
+
+### 核心差异：谁负责什么
+
+| 职责 | CUDA | TileLang |
+|------|------|----------|
+| 线程→元素映射 | 你手动写 `idx = blockIdx.x * blockDim.x + threadIdx.x` | 编译器生成 |
+| 循环边界检查 | 你手动写 `i < n` | 编译器生成 |
+| shared memory 管理 | 你手动声明、加载、同步 | 编译器处理 |
+| launch 参数 | 你手动指定 grid/block size | 编译器自动选择 |
+| 内存合并访问 | 你需要确保访问模式正确 | 编译器优化 |
+
+---
+
+## 案例二：reduce_sum 算子
+
+### CUDA 实现
+
+树形归约需要关注 shared memory 管理和线程同步：
+
+```cuda
+template <typename T>
+__global__ void reduce_sum_kernel(T *out, const T *in, int64_t N, int64_t C) {
+    extern __shared__ float smem[];
+    int64_t row_id = blockIdx.x;
+    int tid = threadIdx.x;
+
+    // 1. 每个线程累加自己负责的列
+    float sum = 0.0f;
+    for (int64_t c = tid; c < C; c += blockDim.x) {
+        sum += (float)in[row_id * C + c];
+    }
+    smem[tid] = sum;
+    __syncthreads();
+
+    // 2. 树形归约
+    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
+        if (tid < s) {
+            smem[tid] += smem[tid + s];
+        }
+        __syncthreads();
+    }
+
+    // 3. 输出结果
+    if (tid == 0) {
+        out[row_id] = (T)smem[0];
+    }
+}
+```
+
+**陷阱点：**
+- `__syncthreads()` 必须放在条件分支外
+- shared memory 大小需要作为 launch 参数传入
+- 树形归约的 stride 和边界条件容易写错
+
+### TileLang 实现
+
+```python
+@tilelang.jit
+def reduce_sum_kernel(inp, BLOCK_C: int, dtype):
+    N = T.const("N")
+    C = T.const("C")
+    inp: T.Tensor((N, C), dtype)
+    out = T.empty((N,), dtype)
+
+    for n in T.Parallel(N):
+        acc = T.alloc_fragment((1,), dtype)
+        acc[0] = 0.0
+        for c in T.Serial(C // BLOCK_C):
+            tile = T.copy(inp[n, c * BLOCK_C : (c + 1) * BLOCK_C])
+            acc[0] += T.reduce_sum(tile)
+        out[n] = acc[0]
+
+    return out
+```
+
+**关键差异：**
+- shared memory 归约被 `T.reduce_sum` 替代——编译器负责生成高效的归约代码
+- `T.Serial` 保证分块累加的顺序（因为累加有依赖）
+- 不需要手动管理 shared memory 大小和同步
+
+---
+
+## 选择建议
+
+### 什么时候用 CUDA？
+
+- 需要极致性能，想手动控制每一个优化细节
+- 算子涉及复杂的 shared memory / warp-level 操作
+- 需要嵌入已有 CUDA 生态（CUTLASS、cuBLAS）
+
+### 什么时候用 TileLang？
+
+- 快速原型：写一个可工作的 kernel 远比性能重要
+- 跨硬件：同一份代码可编译到 NVIDIA / MetaX / 甚至未来其他硬件
+- 学习阶段：先理解算法逻辑，不必被线程细节分心
+
+### intro-ops 的建议路径
+
+1. **入门**：先写 TileLang 版本，快速跑通算法逻辑
+2. **进阶**：再写 CUDA 版本，理解线程级别的执行细节
+3. **专精**：对比两个版本，分析编译器生成的差异，手动调优 CUDA

docs/troubleshooting.md

@@ -0,0 +1,163 @@
+# 常见错误与排错指南
+
+## 编译错误
+
+### nvcc 版本不匹配
+
+**症状**：`CMake Error: nvcc not found` 或 `nvcc fatal: Unsupported gpu architecture`
+
+**原因**：CUDA Toolkit 版本与 CMake 预设的架构参数不兼容。
+
+**解决**：
+1. 检查 nvcc 版本：`nvcc --version`
+2. 在 `CMakePresets.json` 中将 `CMAKE_CUDA_ARCHITECTURES` 改为 `"native"`，或显式指定你的 GPU 架构编号
+3. 确保 `nvcc` 在 PATH 中
+
+### CUDA 架构参数错误
+
+**症状**：`cudaErrorNoKernelImageForDevice` 或运行时 kernel launch 失败
+
+**原因**：编译时指定的 GPU 架构（如 `89` 对应 L40）与运行时 GPU 不匹配。
+
+**解决**：
+- 将 `CMAKE_CUDA_ARCHITECTURES` 设为 `"native"` 让 CMake 自动检测
+- 或根据 GPU 型号查表设置正确的架构编号
+
+### CUTLASS 拉取失败
+
+**症状**：`FetchContent failed to download CUTLASS`
+
+**原因**：GitHub 网络不可达或代理问题。
+
+**解决**：
+1. 检查网络：`git ls-remote https://github.com/NVIDIA/cutlass.git`
+2. 配置代理：`git config --global http.proxy http://your-proxy:port`
+3. 或手动下载 CUTLASS v3.7.0 放到 `third_party/cutlass/`
+
+---
+
+## 运行时错误
+
+### 越界访问（grid-stride loop 边界条件）
+
+**症状**：随机数值错误、CUDA `illegal memory access` 错误
+
+**原因**：`idx < N` 条件缺失或写错，导致线程访问超出 tensor 范围的内存。
+
+**解决**：
+```cuda
+// 正确写法：每次循环都检查 idx < N
+for (int idx = blockIdx.x * blockDim.x + threadIdx.x;
+     idx < N;
+     idx += gridDim.x * blockDim.x) {
+    out[idx] = in[idx];  // 安全：idx 始终 < N
+}
+```
+
+### `__syncthreads()` 在条件分支内
+
+**症状**：kernel 在 reduce_sum 或 softmax 中挂起（hang），或结果错误
+
+**原因**：`__syncthreads()` 放在 `if` 分支内。CUDA 要求 block 内所有线程都到达同一个 `__syncthreads()`，如果有线程走 else 分支跳过同步点，整个 block 就会死锁。
+
+**解决**：
+```cuda
+// 错误
+if (threadIdx.x < N) {
+    smem[tid] = val;
+    __syncthreads();  // 部分线程不执行，死锁！
+}
+
+// 正确
+smem[tid] = (threadIdx.x < N) ? val : 0;
+__syncthreads();  // 全部线程都到达
+```
+
+### shared memory 大小不足
+
+**症状**：编译错误 `uses too much shared data`
+
+**原因**：申请的 shared memory 超过 GPU 的物理限制（通常 48KB-164KB/block）。
+
+**解决**：
+- 检查 `extern __shared__` 声明的数组大小
+- 减小 block size 或 shared memory 使用量
+- 分多轮处理
+
+---
+
+## 数值错误
+
+### softmax 未减 max（大数值溢出）
+
+**症状**：softmax 输出全为 NaN 或 inf
+
+**原因**：`exp(x)` 在 `x > 88` 时溢出为 inf，直接除 inf 得到 NaN。
+
+**解决**：先减行最大值再做 exp：
+```cuda
+float max_val = row[0];
+for (int j = 1; j < N; j++) {
+    max_val = fmaxf(max_val, row[j * stride]);
+}
+float sum = 0;
+for (int j = 0; j < N; j++) {
+    sum += expf(row[j * stride] - max_val);
+}
+```
+
+### 浮点精度差异（多后端对比）
+
+**症状**：NVIDIA 和 TileLang 的输出在小数点后几位不一致，测试报 `not close`
+
+**原因**：不同后端使用不同数学库（`exp2`/`log2` vs `exp`/`log`），或规约顺序不同导致浮点累加误差。
+
+**解决**：
+- 检查测试容差设置是否合理（FP16 容差应比 FP32 宽松）
+- 确认两边的算法逻辑一致（如 online vs 三趟 softmax）
+- 规约顺序差异导致的误差在合理范围内（`rtol=1e-3` for FP16）则正常
+
+### 规约顺序影响结果
+
+**症状**：reduce_sum 结果每次运行都略有不同
+
+**原因**：浮点加法不满足结合律，tree reduction 的配对顺序影响结果。
+
+**解决**：这是正常的浮点行为，只要误差在容差范围内即可接受。
+
+---
+
+## 性能问题
+
+### bank conflict
+
+**症状**：reduce_sum kernel 带宽利用率远低于预期（如 < 50%）
+
+**原因**：shared memory 访问时多个线程命中同一 bank。
+
+**解决**：
+- 添加 padding 偏移访问地址
+- 使用 `__shared__ float smem[BLOCK_SIZE + PADDING]` 并错开 bank
+- 将 stride 从 1 改为 2 开始归约
+
+### 线程利用率低
+
+**症状**：benchmark 带宽利用率低，但代码逻辑正确
+
+**原因**：grid/block 尺寸选择不当，GPU SM 上的线程不够填满所有 core。
+
+**解决**：
+- block size 建议 128/256/512（取 32 的倍数）
+- grid size 建议 `(N + block_size - 1) / block_size` 或更大
+- 用 `cudaOccupancyMaxPotentialBlockSize` 自动计算最佳配置
+
+### 全局内存未合并访问
+
+**症状**：copy 或 vector_add 带宽利用率远低于峰值
+
+**原因**：线程访问的内存地址不连续，导致多次内存事务而非一次合并访问。
+
+**解决**：
+- 确保相邻线程访问相邻内存地址（lane 0→地址 0，lane 1→地址 1…）
+- 检查 stride 是否为 1
+- 检查数据类型是否与访问模式对齐

notebooks/00_environment_setup.ipynb

@@ -0,0 +1,164 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 00 — 环境配置与验证\n",
+    "\n",
+    "本 Notebook 帮助你验证 intro-ops 开发环境是否正确配置。\n",
+    "\n",
+    "## 检查清单\n",
+    "\n",
+    "1. GPU 是否可用\n",
+    "2. CUDA Toolkit 版本\n",
+    "3. PyTorch 是否正确安装\n",
+    "4. intro-ops 是否成功构建\n",
+    "5. 运行一个简单的基线测试"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 1: 检查 GPU 和 CUDA"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "\n",
+    "print(f\"PyTorch 版本: {torch.__version__}\")\n",
+    "print(f\"CUDA 可用: {torch.cuda.is_available()}\")\n",
+    "\n",
+    "if torch.cuda.is_available():\n",
+    "    print(f\"GPU 型号: {torch.cuda.get_device_name(0)}\")\n",
+    "    print(f\"CUDA 版本: {torch.version.cuda}\")\n",
+    "    print(f\"GPU 数量: {torch.cuda.device_count()}\")\n",
+    "    print(f\"当前设备: {torch.cuda.current_device()}\")\n",
+    "else:\n",
+    "    print(\"⚠️  未检测到 CUDA GPU。intro-ops 需要 NVIDIA GPU 才能运行。\")\n",
+    "    print(\"   请确认 CUDA Toolkit 和 NVIDIA 驱动已正确安装。\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 2: 检查 intro-ops 构建产物"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import os\n",
+    "import sys\n",
+    "\n",
+    "# 检查 CAMP_BUILD_DIR 环境变量\n",
+    "build_dir = os.environ.get(\"CAMP_BUILD_DIR\", \"build-nvidia\")\n",
+    "lib_path = os.path.join(build_dir, \"libcamp_ops.so\")\n",
+    "\n",
+    "print(f\"构建目录: {build_dir}\")\n",
+    "print(f\"库文件路径: {lib_path}\")\n",
+    "print(f\"库文件存在: {os.path.exists(lib_path)}\")\n",
+    "\n",
+    "if not os.path.exists(lib_path):\n",
+    "    print(\"\\n⚠️  libcamp_ops.so 未找到。请先构建项目：\")\n",
+    "    print(\"   bash scripts/build_nvidia.sh configure\")\n",
+    "    print(\"   bash scripts/build_nvidia.sh build\")\n",
+    "    print(f\"\\n   然后设置环境变量: export CAMP_BUILD_DIR={build_dir}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 3: 验证 Python 依赖"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import importlib\n",
+    "\n",
+    "deps = [\"torch\", \"pytest\", \"cmake\", \"ninja\"]\n",
+    "for dep in deps:\n",
+    "    try:\n",
+    "        mod = importlib.import_module(dep)\n",
+    "        version = getattr(mod, \"__version__\", \"unknown\")\n",
+    "        print(f\"✓ {dep}: {version}\")\n",
+    "    except ImportError:\n",
+    "        print(f\"✗ {dep}: 未安装\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 4: 运行基线测试\n",
+    "\n",
+    "如果以上检查全部通过，运行一个简单的 copy 算子测试验证端到端流程。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import subprocess\n",
+    "import sys\n",
+    "\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_copy.py\", \"-v\",\n",
+    "    \"--backend\", \"nvidia\",\n",
+    "    \"-x\",  # 遇到第一个失败就停止\n",
+    "], capture_output=True, text=True)\n",
+    "\n",
+    "print(result.stdout)\n",
+    "if result.returncode != 0:\n",
+    "    print(result.stderr)\n",
+    "    print(\"\\n⚠️  基线测试失败。请检查上面的错误信息。\")\n",
+    "else:\n",
+    "    print(\"\\n✓ 环境配置正确，可以开始学习了！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 常见问题\n",
+    "\n",
+    "- **`torch.cuda.is_available()` 返回 False**：检查 NVIDIA 驱动和 CUDA Toolkit 安装\n",
+    "- **`libcamp_ops.so` 不存在**：运行 `bash scripts/build_nvidia.sh configure && bash scripts/build_nvidia.sh build`\n",
+    "- **ImportError**：运行 `pip install -r requirements.txt`\n",
+    "\n",
+    "更多帮助见 `docs/FAQ.md` 和 `docs/troubleshooting.md`。"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}

notebooks/01_copy_kernel.ipynb

@@ -0,0 +1,245 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 01 — Copy 算子\n",
+    "\n",
+    "## 学习目标\n",
+    "\n",
+    "1. 理解 **grid-stride loop**：为什么它可以处理任意大小的 tensor\n",
+    "2. 完成 NVIDIA CUDA kernel 的 TODO\n",
+    "3. 完成 TileLang kernel 的 TODO\n",
+    "4. 验证正确性 + 查看 benchmark 结果\n",
+    "\n",
+    "Copy 是最简单的算子——把数据从 `src` 搬到 `dst`，不涉及任何计算。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 概念导入：Grid-Stride Loop\n",
+    "\n",
+    "GPU 有成千上万个线程，但 tensor 大小可能更大。Grid-stride loop 让每个线程处理多个元素：\n",
+    "\n",
+    "```python\n",
+    "# 伪代码\n",
+    "for i in range(thread_idx, N, grid_total_threads):\n",
+    "    dst[i] = src[i]\n",
+    "```\n",
+    "\n",
+    "- `grid_total_threads = gridDim.x * blockDim.x`（所有线程总数）\n",
+    "- 步长 = grid_total_threads，每个线程“跨步”处理\n",
+    "- 循环条件 `i < N` 保证不越界\n",
+    "\n",
+    "详细图示见 [docs/diagrams/grid-stride-loop.md](../docs/diagrams/grid-stride-loop.md)。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## PyTorch 参考实现\n",
+    "\n",
+    "先看看用 PyTorch 怎么做 copy——训练营的目标就是实现和它一样的功能："
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "\n",
+    "# PyTorch 版 copy\n",
+    "def pytorch_copy(src: torch.Tensor) -> torch.Tensor:\n",
+    "    return src.clone()\n",
+    "\n",
+    "# 测试\n",
+    "src = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
+    "expected = pytorch_copy(src)\n",
+    "print(f\"输入: {src[:5]}\")\n",
+    "print(f\"输出: {expected[:5]}\")\n",
+    "print(f\"一致: {torch.allclose(src, expected)}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## NVIDIA CUDA Kernel\n",
+    "\n",
+    "打开 `ops/copy/nvidia/kernel.cuh`，你会看到这个骨架。请填写 `TODO` 部分："
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 以下代码来自 ops/copy/nvidia/kernel.cuh，仅供参考——请直接编辑该文件\n",
+    "\n",
+    "KERNEL_SKELETON = \"\"\"\n",
+    "template <typename T>\n",
+    "__global__ void copy_contiguous_kernel(T *dst, const T *src, int64_t n) {\n",
+    "    // TODO: implement a grid-stride loop copy kernel.\n",
+    "    //\n",
+    "    // Suggested steps:\n",
+    "    // 1. Compute the global thread index.\n",
+    "    //    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;\n",
+    "    // 2. Compute the grid-wide stride.\n",
+    "    //    int64_t stride = gridDim.x * blockDim.x;\n",
+    "    // 3. Loop over i = idx; i < n; i += stride.\n",
+    "    // 4. Copy src[i] to dst[i].\n",
+    "}\n",
+    "\"\"\"\n",
+    "print(KERNEL_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### 你的任务\n",
+    "\n",
+    "编辑 `ops/copy/nvidia/kernel.cuh`，完成 TODO 后，重新构建并运行下面的验证："
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import subprocess, sys\n",
+    "\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_copy.py\", \"-v\",\n",
+    "    \"--backend\", \"nvidia\"\n",
+    "], capture_output=True, text=True)\n",
+    "\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ NVIDIA copy kernel 测试全部通过！\")\n",
+    "else:\n",
+    "    print(result.stderr[-500:] if result.stderr else \"\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## TileLang Kernel\n",
+    "\n",
+    "打开 `ops/copy/tilelang/kernel.py`，填写 TODO："
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "TILELANG_SKELETON = \"\"\"\n",
+    "import tilelang\n",
+    "import tilelang.language as T\n",
+    "\n",
+    "@tilelang.jit\n",
+    "def copy_kernel(src, BLOCK_N: int, dtype):\n",
+    "    N = T.const(\"N\")\n",
+    "    src: T.Tensor((N,), dtype)\n",
+    "    out = T.empty((N,), dtype)\n",
+    "\n",
+    "    # TODO: implement a tile-wise copy kernel.\n",
+    "    #\n",
+    "    # Suggested steps:\n",
+    "    # 1. Launch one TileLang kernel over the N // BLOCK_N tiles.\n",
+    "    #    for i in T.Parallel(N // BLOCK_N):\n",
+    "    # 2. Use T.copy to move one tile from src to out.\n",
+    "    #    tile = T.copy(src[i * BLOCK_N : (i + 1) * BLOCK_N])\n",
+    "    #    out[i * BLOCK_N : (i + 1) * BLOCK_N] = tile\n",
+    "\n",
+    "    return out\n",
+    "\"\"\"\n",
+    "print(TILELANG_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 TileLang 版本\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_copy.py\", \"-v\",\n",
+    "    \"--backend\", \"tilelang\"\n",
+    "], capture_output=True, text=True)\n",
+    "\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ TileLang copy kernel 测试全部通过！\")\n",
+    "else:\n",
+    "    print(result.stderr[-500:] if result.stderr else \"\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 对比：你的实现 vs PyTorch"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "from operator_runtime import copy\n",
+    "\n",
+    "src = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
+    "my_out = copy(src, backend=\"nvidia\")\n",
+    "torch_out = src.clone()\n",
+    "\n",
+    "print(f\"输入:        {src[:5]}\")\n",
+    "print(f\"我的 kernel:  {my_out[:5]}\")\n",
+    "print(f\"PyTorch:      {torch_out[:5]}\")\n",
+    "print(f\"一致: {torch.allclose(my_out, torch_out)}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 检查清单\n",
+    "\n",
+    "- [ ] 理解 grid-stride loop 原理\n",
+    "- [ ] 理解 global memory 合并访问（coalesced access）\n",
+    "- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
+    "- [ ] TileLang `kernel.py` TODO 完成\n",
+    "- [ ] 两种后端测试全部通过\n",
+    "- [ ] benchmark 跑通，带宽数据合理"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}

notebooks/02_vector_add_kernel.ipynb

@@ -0,0 +1,239 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 02 — Vector Add 算子\n",
+    "\n",
+    "## 学习目标\n",
+    "\n",
+    "1. 理解逐元素并行：每个线程独立处理一对 `(a[i], b[i])`\n",
+    "2. 理解 tile-level 并行（`T.Parallel` vs `T.Serial` 的区别）\n",
+    "3. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
+    "\n",
+    "Vector add 和 copy 结构相似，但引入了**双输入**和**逐元素计算**。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 概念导入：逐元素并行\n",
+    "\n",
+    "```\n",
+    "a = [a0, a1, a2, ..., aN-1]\n",
+    "b = [b0, b1, b2, ..., bN-1]\n",
+    "         ↓ 逐元素相加\n",
+    "c = [a0+b0, a1+b1, a2+b2, ..., aN-1+bN-1]\n",
+    "```\n",
+    "\n",
+    "每个 `c[i] = a[i] + b[i]` 完全独立——不需要线程间通信。\n",
+    "Grid-stride loop 同样适用。\n",
+    "\n",
+    "详细图示见 [docs/diagrams/thread-grid-layout.md](../docs/diagrams/thread-grid-layout.md)。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## PyTorch 参考实现"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "\n",
+    "a = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
+    "b = torch.randn_like(a)\n",
+    "expected = a + b\n",
+    "\n",
+    "print(f\"a[:5]:   {a[:5]}\")\n",
+    "print(f\"b[:5]:   {b[:5]}\")\n",
+    "print(f\"a+b[:5]: {expected[:5]}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## NVIDIA CUDA Kernel\n",
+    "\n",
+    "打开 `ops/vector_add/nvidia/kernel.cuh`，你需要完成两部分：\n",
+    "\n",
+    "1. **`add_values<T>` 辅助函数**：返回 `a + b`（泛型版本和 half 特化版）\n",
+    "2. **`vector_add_contiguous_kernel`**：grid-stride loop 遍历，`out[i] = add_values(a[i], b[i])`"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "KERNEL_SKELETON = \"\"\"\n",
+    "template <typename T>\n",
+    "__device__ T add_values(T a, T b) {\n",
+    "    // TODO: return the elementwise sum for generic types.\n",
+    "}\n",
+    "\n",
+    "template <>\n",
+    "__device__ inline half add_values<half>(half a, half b) {\n",
+    "    // TODO: return the half-precision elementwise sum.\n",
+    "    // 提示：half 的加法用 __hadd(a, b)\n",
+    "}\n",
+    "\n",
+    "template <typename T>\n",
+    "__global__ void vector_add_contiguous_kernel(\n",
+    "    T *out, const T *a, const T *b, int64_t n) {\n",
+    "    // TODO: grid-stride loop, out[i] = add_values(a[i], b[i])\n",
+    "}\n",
+    "\"\"\"\n",
+    "print(KERNEL_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 NVIDIA 版本\n",
+    "import subprocess, sys\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_vector_add.py\", \"-v\",\n",
+    "    \"--backend\", \"nvidia\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ NVIDIA vector_add kernel 测试全部通过！\")\n",
+    "else:\n",
+    "    print(result.stderr[-500:] if result.stderr else \"\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## TileLang Kernel\n",
+    "\n",
+    "打开 `ops/vector_add/tilelang/kernel.py`，注意这里使用了 `T.Parallel`。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "TILELANG_SKELETON = \"\"\"\n",
+    "@tilelang.jit\n",
+    "def vector_add_kernel(a, b, BLOCK_N: int, dtype):\n",
+    "    N = T.const(\"N\")\n",
+    "    a: T.Tensor((N,), dtype)\n",
+    "    b: T.Tensor((N,), dtype)\n",
+    "    out = T.empty((N,), dtype)\n",
+    "\n",
+    "    # TODO: implement a tile-wise vector add kernel.\n",
+    "    #\n",
+    "    # Suggested steps:\n",
+    "    # 1. T.Parallel(N // BLOCK_N) to iterate over tiles\n",
+    "    # 2. Compute base = tile_idx * BLOCK_N\n",
+    "    # 3. Inner T.Parallel(BLOCK_N): out[base+i] = a[base+i] + b[base+i]\n",
+    "\n",
+    "    return out\n",
+    "\"\"\"\n",
+    "print(TILELANG_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### `T.Parallel` vs `T.Serial`\n",
+    "\n",
+    "| | `T.Parallel` | `T.Serial` |\n",
+    "|--|-------------|------------|\n",
+    "| 含义 | 循环迭代可以并行 | 循环迭代必须顺序执行 |\n",
+    "| 何时用 | 迭代之间无数据依赖 | 迭代之间有数据依赖（如累加器） |\n",
+    "| vector_add | ✓ 每个 `(a[i], b[i])` 独立 | N/A |"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 TileLang 版本\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_vector_add.py\", \"-v\",\n",
+    "    \"--backend\", \"tilelang\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ TileLang vector_add kernel 测试全部通过！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 对比验证"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "from operator_runtime import vector_add\n",
+    "\n",
+    "a = torch.randn(1024, device=\"cuda\", dtype=torch.float32)\n",
+    "b = torch.randn_like(a)\n",
+    "\n",
+    "my_out = vector_add(a, b, backend=\"nvidia\")\n",
+    "torch_out = a + b\n",
+    "\n",
+    "print(f\"我的 kernel: {my_out[:5]}\")\n",
+    "print(f\"PyTorch:     {torch_out[:5]}\")\n",
+    "print(f\"一致: {torch.allclose(my_out, torch_out)}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 检查清单\n",
+    "\n",
+    "- [ ] 理解逐元素并行的线程网格布局\n",
+    "- [ ] 理解 `T.Parallel` vs `T.Serial` 的区别\n",
+    "- [ ] NVIDIA `kernel.cuh` TODO 完成（add_values + kernel 函数）\n",
+    "- [ ] TileLang `kernel.py` TODO 完成\n",
+    "- [ ] 两种后端测试全部通过\n",
+    "- [ ] benchmark 跑通"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}

notebooks/03_reduce_sum_kernel.ipynb

@@ -0,0 +1,273 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 03 — Reduce Sum 算子\n",
+    "\n",
+    "## 学习目标\n",
+    "\n",
+    "1. 理解 shared memory 树形归约原理\n",
+    "2. 理解 `__syncthreads()` 的使用时机和条件分支限制\n",
+    "3. 理解 `T.Serial` 在归约场景中的作用\n",
+    "4. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
+    "\n",
+    "Reduce sum 是第一个需要**线程间通信**的算子——从 copy/vector_add 的“各自为战“跨越到“协同计算“。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 概念导入：树形归约\n",
+    "\n",
+    "目标：把一行 N 个元素求和为 1 个值。每个 block 处理一行：\n",
+    "\n",
+    "```\n",
+    "Step 0: [3, 1, 7, 0, 4, 1, 6, 3]  ← 线程各自的部分和存入 shared memory\n",
+    "Step 1: [7, 2, 13, 3]  ← stride=4, 相邻配对相加\n",
+    "Step 2: [20, 5]         ← stride=2\n",
+    "Step 3: [25]            ← stride=1, 最终结果\n",
+    "```\n",
+    "\n",
+    "每步后必须 `__syncthreads()` ——因为下一步要读上一步别人写的数据。\n",
+    "\n",
+    "详细图示见 [docs/diagrams/tree-reduction.md](../docs/diagrams/tree-reduction.md)。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## PyTorch 参考实现"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "\n",
+    "src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
+    "expected = torch.sum(src, dim=1)  # 对每行求和\n",
+    "\n",
+    "print(f\"输入 shape: {src.shape}\")\n",
+    "print(f\"输出 shape: {expected.shape}\")\n",
+    "print(f\"第一行: src[0, :5] = {src[0, :5]}\")\n",
+    "print(f\"第一行和: {expected[0]:.6f}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## NVIDIA CUDA Kernel\n",
+    "\n",
+    "打开 `ops/reduce_sum/nvidia/kernel.cuh`。这是第一个使用 shared memory 的 kernel。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "KERNEL_SKELETON = \"\"\"\n",
+    "__global__ void reduce_sum_rowwise_kernel(\n",
+    "    float *out, const float *in, int64_t rows, int64_t cols) {\n",
+    "\n",
+    "    // TODO: implement a row-wise reduce_sum kernel with shared memory.\n",
+    "    //\n",
+    "    // Suggested steps:\n",
+    "    // 1. Use one block per row.\n",
+    "    //    int row = blockIdx.x;\n",
+    "    // 2. Each thread accumulates its columns:\n",
+    "    //    float sum = 0;\n",
+    "    //    for (int c = threadIdx.x; c < cols; c += blockDim.x)\n",
+    "    //        sum += in[row * cols + c];\n",
+    "    // 3. Store partial sum to shared memory, then __syncthreads().\n",
+    "    // 4. Tree reduction:\n",
+    "    //    for (int s = blockDim.x/2; s > 0; s >>= 1) {\n",
+    "    //        if (threadIdx.x < s) smem[tid] += smem[tid + s];\n",
+    "    //        __syncthreads();\n",
+    "    //    }\n",
+    "    // 5. Thread 0 writes smem[0] to out[row].\n",
+    "}\n",
+    "\"\"\"\n",
+    "print(KERNEL_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### ⚠️ 关键陷阱：`__syncthreads()` 不能在条件分支内\n",
+    "\n",
+    "```cuda\n",
+    "// 错误——死锁！\n",
+    "if (tid < s) {\n",
+    "    smem[tid] += smem[tid + s];\n",
+    "    __syncthreads();  // 只有 tid < s 的线程执行同步\n",
+    "}\n",
+    "\n",
+    "// 正确——所有线程都到达同步点\n",
+    "if (tid < s) {\n",
+    "    smem[tid] += smem[tid + s];\n",
+    "}\n",
+    "__syncthreads();  // 全部线程到位\n",
+    "```"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 NVIDIA 版本\n",
+    "import subprocess, sys\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_reduce_sum.py\", \"-v\",\n",
+    "    \"--backend\", \"nvidia\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ NVIDIA reduce_sum kernel 测试全部通过！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## TileLang Kernel\n",
+    "\n",
+    "打开 `ops/reduce_sum/tilelang/kernel.py`。TileLang 版本用 `T.Serial` + `T.reduce_sum` 替代手动树形归约。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "TILELANG_SKELETON = \"\"\"\n",
+    "@tilelang.jit\n",
+    "def reduce_sum_kernel(src, BLOCK_N: int, BLOCK_M: int):\n",
+    "    N, M = T.const(\"N, M\")\n",
+    "    dtype = T.float32\n",
+    "    src: T.Tensor((N, M), dtype)\n",
+    "    out = T.empty((N,), dtype)\n",
+    "\n",
+    "    # TODO: implement a tiled row-wise reduce_sum kernel.\n",
+    "    #\n",
+    "    # Key insight: outer loop uses T.Serial (accumulate state),\n",
+    "    # inner loop uses T.reduce_sum (parallel reduction within tile).\n",
+    "\n",
+    "    return out\n",
+    "\"\"\"\n",
+    "print(TILELANG_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### 为什么外层用 `T.Serial`？\n",
+    "\n",
+    "分块累加时，每个 chunk 的求和结果需要**累加到同一个累加器**。这是有状态依赖的——后一个 chunk 必须在前一个完成之后才能累加。所以列方向的循环是 `T.Serial`。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 TileLang 版本\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_reduce_sum.py\", \"-v\",\n",
+    "    \"--backend\", \"tilelang\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ TileLang reduce_sum kernel 测试全部通过！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 对比验证"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "from operator_runtime import reduce_sum\n",
+    "\n",
+    "src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
+    "my_out = reduce_sum(src, dim=1, backend=\"nvidia\")\n",
+    "torch_out = torch.sum(src, dim=1)\n",
+    "\n",
+    "print(f\"我的 kernel[:5]: {my_out[:5]}\")\n",
+    "print(f\"PyTorch[:5]:     {torch_out[:5]}\")\n",
+    "print(f\"最大误差: {(my_out - torch_out).abs().max().item():.2e}\")\n",
+    "print(f\"一致: {torch.allclose(my_out, torch_out, atol=1e-5, rtol=1e-5)}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 进阶：Bank Conflict 优化\n",
+    "\n",
+    "如果你已经跑通基础版本，可以尝试优化 bank conflict：\n",
+    "\n",
+    "```cuda\n",
+    "// 加 padding 错开 bank 访问\n",
+    "__shared__ float smem[BLOCK_SIZE + PADDING];\n",
+    "// 或者归约从 stride=2 开始而非 stride=1\n",
+    "```\n",
+    "\n",
+    "对比优化前后的 benchmark 带宽数据。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 检查清单\n",
+    "\n",
+    "- [ ] 理解 shared memory 树形归约原理\n",
+    "- [ ] 理解 `__syncthreads()` 的使用时机和条件分支限制\n",
+    "- [ ] 理解 `T.Serial` 在归约场景中的作用\n",
+    "- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
+    "- [ ] TileLang `kernel.py` TODO 完成\n",
+    "- [ ] 两种后端测试全部通过\n",
+    "- [ ] 尝试优化 bank conflict\n",
+    "- [ ] benchmark 跑通"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}

notebooks/04_softmax_kernel.ipynb

@@ -0,0 +1,260 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 04 — Softmax 算子\n",
+    "\n",
+    "## 学习目标\n",
+    "\n",
+    "1. 理解数值稳定性：为什么需要减 max\n",
+    "2. 理解 online softmax 算法（一遍扫描 vs 三趟扫描）\n",
+    "3. 理解 TileLang 中为什么用 `exp2`/`log2` 替代 `exp`/`log`\n",
+    "4. 完成 NVIDIA + TileLang 两个 kernel 的 TODO\n",
+    "\n",
+    "Softmax 是四个算子中最复杂的——既需要线程间通信（同步归约），又有数值稳定性陷阱。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 概念导入：Softmax 的数值稳定性\n",
+    "\n",
+    "### 朴素公式\n",
+    "\n",
+    "$$\\text{softmax}(x_i) = \\frac{e^{x_i}}{\\sum_j e^{x_j}}$$\n",
+    "\n",
+    "### 问题\n",
+    "\n",
+    "$e^{88.7} \\approx 1.6 \\times 10^{38}$，接近 FP32 上限。如果 $x_i = 100$，则 $e^{100}$ 溢出为 inf，inf/inf = NaN。\n",
+    "\n",
+    "### 解决：减最大值\n",
+    "\n",
+    "$$\\text{softmax}(x_i) = \\frac{e^{x_i - \\max(x)}}{\\sum_j e^{x_j - \\max(x)}}$$\n",
+    "\n",
+    "分子分母同除 $e^{\\max(x)}$，数学结果不变，但 $e^{x_i - \\max(x)}$ 最大为 $e^0 = 1$，永不溢出。\n",
+    "\n",
+    "详细图示见 [docs/diagrams/softmax-pipeline.md](../docs/diagrams/softmax-pipeline.md)。"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## PyTorch 参考实现"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "\n",
+    "src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
+    "expected = torch.softmax(src, dim=1)\n",
+    "\n",
+    "print(f\"输入 shape: {src.shape}\")\n",
+    "print(f\"输出 shape: {expected.shape}\")\n",
+    "print(f\"每行和: {expected.sum(dim=1)[:5]}\")  # 应该全为 1.0\n",
+    "\n",
+    "# 演示溢出问题\n",
+    "big = torch.tensor([100.0, 200.0, 300.0], device=\"cuda\")\n",
+    "naive = torch.exp(big) / torch.exp(big).sum()\n",
+    "print(f\"\\n朴素 softmax([100, 200, 300]): {naive}\")  # NaN!\n",
+    "\n",
+    "stable = torch.softmax(big, dim=0)\n",
+    "print(f\"稳定 softmax([100, 200, 300]): {stable}\")    # [0, 0, 1]"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## NVIDIA CUDA Kernel（三趟扫描）\n",
+    "\n",
+    "打开 `ops/softmax/nvidia/kernel.cuh`。基础版本使用三趟扫描：\n",
+    "\n",
+    "1. Pass 1: 求行最大值 `max_val`\n",
+    "2. Pass 2: 写 `exp(x - max)` 到 out，同时累加 sum\n",
+    "3. Pass 3: 除 sum 归一化"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "KERNEL_SKELETON = \"\"\"\n",
+    "__global__ void softmax_rowwise_kernel(\n",
+    "    float *out, const float *in, int64_t rows, int64_t cols) {\n",
+    "\n",
+    "    // TODO: implement a numerically stable row-wise softmax kernel.\n",
+    "    //\n",
+    "    // Suggested steps:\n",
+    "    // 1. One block per row.\n",
+    "    // 2. Pass 1: find row max via shared memory reduction.\n",
+    "    // 3. Pass 2: exp(x - max) → out, accumulate sum via reduction.\n",
+    "    // 4. Pass 3: out[i] /= sum.\n",
+    "}\n",
+    "\"\"\"\n",
+    "print(KERNEL_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### 为什么是三趟？\n",
+    "\n",
+    "因为 softmax 有两个需要全局信息才能计算的步骤：\n",
+    "- 求 max 需要全局信息（所有元素的最大值）\n",
+    "- 求 sum 需要全局信息（所有 exp(x-max) 的和）\n",
+    "\n",
+    "这两个全局信息都通过 shared memory 归约获得，每一步都需要 block 内同步。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 NVIDIA 版本\n",
+    "import subprocess, sys\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_softmax.py\", \"-v\",\n",
+    "    \"--backend\", \"nvidia\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ NVIDIA softmax kernel 测试全部通过！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## TileLang Kernel（Online Softmax）\n",
+    "\n",
+    "打开 `ops/softmax/tilelang/kernel.py`。TileLang 版本使用 **online softmax**——只需两趟扫描。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "TILELANG_SKELETON = \"\"\"\n",
+    "@tilelang.jit\n",
+    "def softmax_kernel(src, BLOCK_N: int, BLOCK_M: int):\n",
+    "    log2_e = 1.44269504    # log2(e) for exp/log2 conversion\n",
+    "    N, M = T.const(\"N, M\")\n",
+    "    dtype = T.float32\n",
+    "    src: T.Tensor((N, M), dtype)\n",
+    "    out = T.empty((N, M), dtype)\n",
+    "\n",
+    "    # TODO: implement a tiled row-wise softmax with online algorithm.\n",
+    "    #\n",
+    "    # Two-pass approach:\n",
+    "    # Pass 1: scroll through column tiles, update running log-sum-exp.\n",
+    "    # Pass 2: scroll again, normalize each tile with final lse.\n",
+    "\n",
+    "    return out\n",
+    "\"\"\"\n",
+    "print(TILELANG_SKELETON)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "### 为什么用 `exp2`/`log2`？\n",
+    "\n",
+    "GPU 硬件对 2 的幂运算 (`2^x`) 有专门的快速指令，比自然指数 (`e^x`) 快。\n",
+    "\n",
+    "转换公式：$\\exp(x) = 2^{x \\cdot \\log_2(e)}$\n",
+    "\n",
+    "其中 $\\log_2(e) \\approx 1.44269504$（代码中的 `log2_e` 常量）。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "# 验证 TileLang 版本\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"-m\", \"pytest\",\n",
+    "    \"tests/op_tests/test_softmax.py\", \"-v\",\n",
+    "    \"--backend\", \"tilelang\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode == 0:\n",
+    "    print(\"✓ TileLang softmax kernel 测试全部通过！\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 对比验证"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "from operator_runtime import softmax\n",
+    "\n",
+    "src = torch.randn(32, 128, device=\"cuda\", dtype=torch.float32)\n",
+    "my_out = softmax(src, dim=1, backend=\"nvidia\")\n",
+    "torch_out = torch.softmax(src, dim=1)\n",
+    "\n",
+    "print(f\"我的 kernel 行和: {my_out.sum(dim=1)[:5]}\")\n",
+    "print(f\"PyTorch 行和:     {torch_out.sum(dim=1)[:5]}\")\n",
+    "print(f\"最大误差: {(my_out - torch_out).abs().max().item():.2e}\")\n",
+    "print(f\"一致: {torch.allclose(my_out, torch_out, atol=1e-5, rtol=1e-5)}\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 检查清单\n",
+    "\n",
+    "- [ ] 理解数值稳定性问题：为什么需要减 max\n",
+    "- [ ] 理解 online softmax 算法（一遍扫描 vs 三趟扫描）\n",
+    "- [ ] 理解 log-sum-exp 的滚动更新逻辑\n",
+    "- [ ] 理解 TileLang 中为什么用 `exp2`/`log2` 替代 `exp`/`log`\n",
+    "- [ ] NVIDIA `kernel.cuh` TODO 完成\n",
+    "- [ ] TileLang `kernel.py` TODO 完成\n",
+    "- [ ] 两种后端测试全部通过\n",
+    "- [ ] 性能与 PyTorch 参考实现对比\n",
+    "- [ ] 尝试 warp-level 优化"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}

notebooks/05_benchmark_and_profile.ipynb

@@ -0,0 +1,213 @@
+{
+ "cells": [
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "# 05 — Benchmark 与 Profiling 入门\n",
+    "\n",
+    "## 学习目标\n",
+    "\n",
+    "1. 运行 benchmark 获取 kernel 性能数据\n",
+    "2. 理解关键指标：带宽利用率、延迟\n",
+    "3. 学会对比自己的实现与 PyTorch 的性能差异\n",
+    "4. 初识 profiling 工具（nsys / ncu）"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 1: 运行 Benchmark\n",
+    "\n",
+    "intro-ops 的 `tests/bench/` 目录下每个算子都有 benchmark 脚本。"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import subprocess, sys\n",
+    "\n",
+    "# 对 copy 算子同时跑正确性和 benchmark\n",
+    "result = subprocess.run([\n",
+    "    sys.executable, \"tests/run_ops.py\",\n",
+    "    \"--op\", \"copy\",\n",
+    "    \"--backend\", \"nvidia\",\n",
+    "    \"--mode\", \"all\"\n",
+    "], capture_output=True, text=True)\n",
+    "print(result.stdout)\n",
+    "if result.returncode != 0:\n",
+    "    print(result.stderr[-500:] if result.stderr else \"\")"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 2: 理解 Benchmark 指标\n",
+    "\n",
+    "### 带宽利用率 (Bandwidth Utilization)\n",
+    "\n",
+    "$$\\text{利用率} = \\frac{\\text{实际带宽}}{\\text{峰值带宽}} \\times 100\\%$$\n",
+    "\n",
+    "- **实际带宽** = 数据量 / 执行时间（单位 GB/s）\n",
+    "- **峰值带宽** = GPU 理论最大内存带宽（如 L40 = 864 GB/s）\n",
+    "- copy 是纯内存搬运，瓶颈在带宽 → 追求高带宽利用率\n",
+    "- reduce_sum/softmax 有计算量，瓶颈可能在计算 → 对比算术强度\n",
+    "\n",
+    "### 关键指标\n",
+    "\n",
+    "| 指标 | 含义 | 好/坏参考 |\n",
+    "|------|------|----------|\n",
+    "| Bandwidth (GB/s) | 实际内存带宽 | 越高越好 |\n",
+    "| Bandwidth Utilization | 带宽利用率 | >80% 优秀，<50% 需优化 |\n",
+    "| Latency (ms) | 单次调用耗时 | 越低越好 |\n",
+    "| GFLOPS | 计算吞吐 | 越高越好 |"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 3: 对比所有算子"
+   ]
+  },
+  {
+   "cell_type": "code",
+   "execution_count": null,
+   "metadata": {},
+   "outputs": [],
+   "source": [
+    "import torch\n",
+    "from operator_runtime import copy, vector_add, reduce_sum, softmax\n",
+    "import time\n",
+    "\n",
+    "def bench_op(name, fn, *args, warmup=10, iters=100):\n",
+    "    for _ in range(warmup):\n",
+    "        fn(*args)\n",
+    "    torch.cuda.synchronize()\n",
+    "    start = time.perf_counter()\n",
+    "    for _ in range(iters):\n",
+    "        fn(*args)\n",
+    "    torch.cuda.synchronize()\n",
+    "    elapsed_ms = (time.perf_counter() - start) / iters * 1000\n",
+    "    print(f\"{name:15s}: {elapsed_ms:8.3f} ms\")\n",
+    "    return elapsed_ms\n",
+    "\n",
+    "# FP32 测试\n",
+    "print(\"=== FP32 Benchmarks ===\")\n",
+    "src = torch.randn(1024 * 1024, device=\"cuda\", dtype=torch.float32)\n",
+    "bench_op(\"copy\", lambda x: copy(x, backend=\"nvidia\"), src)\n",
+    "\n",
+    "a = torch.randn(1024 * 1024, device=\"cuda\", dtype=torch.float32)\n",
+    "b = torch.randn_like(a)\n",
+    "bench_op(\"vector_add\", lambda x, y: vector_add(x, y, backend=\"nvidia\"), a, b)\n",
+    "\n",
+    "mat = torch.randn(1024, 1024, device=\"cuda\", dtype=torch.float32)\n",
+    "bench_op(\"reduce_sum\", lambda x: reduce_sum(x, dim=1, backend=\"nvidia\"), mat)\n",
+    "\n",
+    "mat2 = torch.randn(1024, 1024, device=\"cuda\", dtype=torch.float32)\n",
+    "bench_op(\"softmax\", lambda x: softmax(x, dim=1, backend=\"nvidia\"), mat2)\n",
+    "\n",
+    "print(\"\\n=== PyTorch Baseline ===\")\n",
+    "bench_op(\"torch.clone\", lambda x: x.clone(), src)\n",
+    "bench_op(\"torch.add\", lambda x, y: x + y, a, b)\n",
+    "bench_op(\"torch.sum\", lambda x: torch.sum(x, dim=1), mat)\n",
+    "bench_op(\"torch.softmax\", lambda x: torch.softmax(x, dim=1), mat2)"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 4: Profiling 入门\n",
+    "\n",
+    "### NVIDIA Nsight Systems (nsys)\n",
+    "\n",
+    "```bash\n",
+    "# 对 copy 算子做 profiling\n",
+    "nsys profile --stats=true \\\n",
+    "  python tests/bench/copy.py --backend nvidia\n",
+    "\n",
+    "# 输出关键指标：\n",
+    "# - CUDA Kernel 执行时间\n",
+    "# - 内存传输时间\n",
+    "# - Occupancy\n",
+    "```\n",
+    "\n",
+    "### NVIDIA Nsight Compute (ncu)\n",
+    "\n",
+    "```bash\n",
+    "# 详细 kernel 分析\n",
+    "ncu --set full \\\n",
+    "  python tests/bench/copy.py --backend nvidia\n",
+    "\n",
+    "# 关注：\n",
+    "# - Memory Throughput\n",
+    "# - Compute Throughput\n",
+    "# - Occupancy\n",
+    "# - Shared Memory Configuration\n",
+    "```"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## Step 5: 性能分析检查清单\n",
+    "\n",
+    "当你的 kernel 正确运行但性能不理想时，按此顺序排查：\n",
+    "\n",
+    "### copy / vector_add\n",
+    "\n",
+    "1. 带宽利用率低于 60%？→ 检查合并访问（coalesced access）\n",
+    "2. grid/block 尺寸是否合理？→ 尝试 block_size=256/512\n",
+    "3. stride 是否为 1？\n",
+    "\n",
+    "### reduce_sum\n",
+    "\n",
+    "1. 带宽利用率低于 50%？→ 可能有 bank conflict\n",
+    "2. 尝试加 padding：`__shared__ float smem[BLOCK_SIZE + 1]`\n",
+    "3. 尝试 warp-level shuffle 替代 shared memory 归约\n",
+    "\n",
+    "### softmax\n",
+    "\n",
+    "1. 三趟扫描的带宽开销 = N × 3 次全局内存读写 → 能否减为两趟？\n",
+    "2. 用 online softmax 减少一次全局内存写\n",
+    "3. 用 `__shfl_down_sync` 在 warp 内做归约，节省 shared memory 带宽"
+   ]
+  },
+  {
+   "cell_type": "markdown",
+   "metadata": {},
+   "source": [
+    "## 参考：L40 GPU 峰值数据\n",
+    "\n",
+    "| 指标 | L40 | A100 |\n",
+    "|------|-----|------|\n",
+    "| 峰值带宽 | 864 GB/s | 2039 GB/s |\n",
+    "| FP32 TFLOPS | 90.5 | 19.5 |\n",
+    "| SM 数量 | 142 | 108 |\n",
+    "| Shared Memory / SM | 128 KB | 164 KB |\n",
+    "\n",
+    "> 你 GPU 的峰值数据见 `tests/perf_profiles/local_gpu.yaml`，benchmark 会自动读取该文件做对比。"
+   ]
+  }
+ ],
+ "metadata": {
+  "kernelspec": {
+   "display_name": "Python 3",
+   "language": "python",
+   "name": "python3"
+  },
+  "language_info": {
+   "name": "python",
+   "version": "3.12.0"
+  }
+ },
+ "nbformat": 4,
+ "nbformat_minor": 5
+}