docs: add memory basics lesson

docs/concepts.md

@@ -129,12 +129,79 @@ Host     │  CPU RAM │  通过 musaMemcpy 搬   ← PCIe 慢得多
 常见误区：
 
 - "**Local memory**" 不是真的快内存,它就是给寄存器装不下的局部变量在 global 上开辟的空间,慢。
-- "**Constant memory**" 只在所有线程**读同一地址**时通过专用 cache 飞快;每个线程读不同地址就退化成 global。
+- "**Constant memory**" 中，同一地址读取可利用广播和 constant cache；读取不同地址会失去广播优势，性能取决于实际访问模式和设备，不能一概等同于 global memory。
 - 优化访存时主要检查三件事：
   1. **Coalesced**:同 warp 相邻线程访问相邻地址(week4)
   2. **Shared 替代 Global**:把多次复用的数据搬进 shared(week5 GEMM)
   3. **避免 Bank Conflict**:shared memory 分 32 个 bank,同 warp 内同 bank 不同地址会串行化
 
+### 全局、共享与常量内存：先判断数据该放哪里
+
+先从数据的读写者和生命周期出发，再考虑优化。下面的对比描述的是作用域与访问方式；具体设备和 SDK 的实现细节会有差异，不能据此推导固定容量、延迟或性能结论。
+
+| 内存 | 谁能访问 | 读写行为 | 适合的数据 | 首先检查什么 |
+| --- | --- | --- | --- | --- |
+| Global（全局内存） | kernel 中的线程 | 可读、可写 | 输入、输出、跨 block 的中间结果 | 相邻线程是否访问相邻的 global 地址 |
+| Shared（共享内存） | 同一 block 的线程 | 可读、可写 | block 内反复使用的 tile、部分和 | 屏障是否安全、是否有 bank conflict |
+| Constant（常量内存） | kernel 中的线程 | 设备端只读，由主机更新 | 小型只读数据，且同一批线程经常读取同一地址 | 是否呈现广播式（同地址）访问 |
+
+选择顺序可以固定为四步：
+
+1. 跨 block 共享或一般可写的中间结果 → global；需要跨 kernel 存活的数据通常也放在 global，除非它是适合 constant 的小型只读常量符号。
+2. 小型、只读、适合广播的参数 → constant。
+3. 同一 block 内会被反复消费的数据 → shared。
+4. 其他情况先从 global 开始，再测量。
+
+下面三个 MUSA/CUDA 风格片段只说明作用域和访问模式；它们不构成可移植的性能结论，也不替代 Week 4、Week 5 中可运行的示例。
+
+```cpp
+__global__ void global_add(float* out, const float* a, const float* b, int n) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n) out[i] = a[i] + b[i];
+}
+```
+
+`global_add` 通过全局索引读写 global 指针指向的元素，并用 `if (i < n)` 保护尾部线程。
+
+```cpp
+__global__ void block_sum(const float* x, float* partial, int n) {
+  __shared__ float s[256];
+  int t = threadIdx.x;
+  if (blockDim.x != 256 || blockDim.y != 1 || blockDim.z != 1) return;
+  int i = blockIdx.x * 256 + t;
+  s[t] = (i < n) ? x[i] : 0.0f;
+  __syncthreads();
+  if (t == 0) {
+    float sum = 0.0f;
+    for (int j = 0; j < 256; ++j) sum += s[j];
+    partial[blockIdx.x] = sum;
+  }
+}
+```
+
+`block_sum` 要求以一维、恰好 256 线程的 block 启动；若 `blockDim.x != 256`、`blockDim.y != 1` 或 `blockDim.z != 1`，所有线程会在访问 `s` 前一致返回。正常路径中，每个线程先写入 shared 数组，越界元素填零，所有线程无条件到达 `__syncthreads()`，再由 thread 0 对恰好 256 个元素求和并写出每个 block 的部分和。
+
+```cpp
+__constant__ float c_scale;
+
+__global__ void scale(float* out, const float* in, int n) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  if (i < n) out[i] = in[i] * c_scale;
+}
+```
+
+`scale` 从设备端只读的 `c_scale` 读取缩放参数；主机负责在启动 kernel 前更新该常量。
+
+读代码时，可以逐项检查：
+
+1. 指针实际驻留在哪一种内存中？
+2. 谁会读取、谁会写入这份数据？
+3. 每个 block 的每个线程是否都能到达屏障？
+4. global 访问是否呈现相邻地址模式？
+5. 尾部线程是否受到边界保护？
+
+继续学习时，可先阅读 [Week 3：同步前提](../code/week3/learning-notes.md)，再结合 [Week 4：global memory 优化](../code/week4/learning-notes.md) 和 [Week 5：shared/constant 与复用](../code/week5/learning-notes.md) 的可运行示例理解这些选择。
+
 ---
 
 ## 5. 执行模型与同步

docs/quiz.html

@@ -712,6 +712,18 @@ <h3>题目列表</h3>
         answer: "MUSA warp size 通常按 128 理解，shuffle mask、warp 数量、reduction tree 深度都要重新算。",
         pitfall: "代码语法接近 CUDA，不代表调优参数可以机械照搬。"
       },
+      {
+        id: "memory-basics-q1",
+        deck: "Week 3",
+        type: "choice",
+        level: "基础",
+        source: "docs/concepts.md",
+        question: "`__syncthreads()` 之后，哪些线程保证能看到彼此对 shared memory 的写入？",
+        choices: ["同一个 block 中到达该 barrier 的线程", "同一个 grid 中所有正在运行的线程", "同一个 SM 上的所有 block", "所有已经启动的 kernel 中的线程"],
+        answerIndex: 0,
+        answer: "`__syncthreads()` 是 block-local 的同步与可见性边界：同一 block 中到达 barrier 的线程可以看到彼此此前对 shared memory 的写入；它不能协调不同 block。",
+        pitfall: "不要把 barrier 放在会让同一 block 线程分歧的分支里，否则部分线程到不了 barrier，可能死锁或产生未定义行为。"
+      },
       {
         id: "w4-q1",
         deck: "Week 4",
@@ -734,6 +746,18 @@ <h3>题目列表</h3>
         answer: "padding 改变 shared memory 行跨度，降低多个线程访问同一 bank 的概率。",
         pitfall: "padding 多一个元素看似浪费一点空间，但常能换来更少 bank conflict。"
       },
+      {
+        id: "memory-basics-q2",
+        deck: "Week 4",
+        type: "choice",
+        level: "基础",
+        source: "docs/concepts.md",
+        question: "`global_add` 已经算出正确结果后，要宣称 global memory 访问高效，还必须检查什么？",
+        choices: ["相邻线程的地址模式，以及测得的有效带宽", "只检查输出数值是否正确", "host CPU 的时钟频率", "调用了多少次 printf"],
+        answerIndex: 0,
+        answer: "正确结果只说明计算逻辑正确；还要检查相邻线程是否访问相邻地址，并用测得的有效带宽判断实际访问表现。",
+        pitfall: "输出数值正确并不能证明访问模式高效；仍要检查线程地址序列和实际带宽。"
+      },
       {
         id: "w5-q1",
         deck: "Week 5",
@@ -761,6 +785,30 @@ <h3>题目列表</h3>
         answer: "第一次等所有线程把 tile 装进 shared；第二次等本轮计算读完 shared，再允许下一轮覆盖 tile。",
         pitfall: "漏掉任意一个同步点都可能产生数据竞态和偶发错误。"
       },
+      {
+        id: "memory-basics-q3",
+        deck: "Week 5",
+        type: "choice",
+        level: "核心",
+        source: "docs/concepts.md",
+        question: "一个 block 执行结束后，关于它写入的 shared memory 内容，哪项说法正确？",
+        choices: ["其他 block 不能读取它，也不能在下一次 kernel 中依赖它；它只是 block 私有的临时存储", "同一 grid 中其他 block 可以直接读取它", "下一次 kernel 会从 global cache 自动恢复它", "它会自动成为所有 kernel 都可写的全局数据"],
+        answerIndex: 0,
+        answer: "shared memory 仅供同一个 block 在该次 kernel 执行中使用。block 结束后，其他 block 不能读取它，后续 kernel 也不能把它当作可依赖的数据来源。",
+        pitfall: "不要把 block 私有的 shared memory 当成跨 block 或跨 kernel 的 global cache。"
+      },
+      {
+        id: "memory-basics-q4",
+        deck: "Week 5",
+        type: "choice",
+        level: "核心",
+        source: "docs/concepts.md",
+        question: "`__constant__` 符号中的数据通常如何更新？",
+        choices: ["主机在 kernel 启动前复制或更新 device symbol；kernel 只读它，不写它", "kernel 通过普通可写 device 指针直接更新它", "任何 block 都可把它当作 shared memory 写入", "它会在每次 kernel 启动时自动从输出数组更新"],
+        answerIndex: 0,
+        answer: "constant symbol 通常由主机在 kernel 启动前复制或更新；kernel 可以读取它，但不能把它当作普通可写的 device 指针。",
+        pitfall: "constant symbol 不是由 kernel 随意写入的普通 device 指针。"
+      },
       {
         id: "w6-q1",
         deck: "Week 6",

docs/superpowers/specs/2026-09-20-memory-basics-design.md

@@ -0,0 +1,55 @@
+# 全局、共享与常量内存基础内容设计
+
+日期：2026-09-20
+
+## 目标
+
+为刚进入 Week 4/5 的学习者，在已有内存层级速查之后补上一段可直接用于读 kernel 的基础教材。读者完成后应能区分 global、shared、constant memory 的作用域与访问方式，并能根据数据复用和访问模式作出初步放置选择。
+
+## 范围与边界
+
+本次只修改概念教材与 Quiz，不新增 `.mu` 可执行文件，不改变 Week 4/5 的实验主线，也不产生性能结论。
+
+- 主教材：在 `docs/concepts.md` 的“内存层级”后补充一个小节。
+- 自测：在 `docs/quiz.html` 的 `QUESTIONS` 中新增四题，题目来源都指向 `docs/concepts.md`。
+- 不写死 shared/constant 容量、bank 数、warp 宽度或周期；这些取决于设备和 SDK。
+- 示例采用 MUSA/CUDA 相近的 kernel 语法，并标为“阅读示意”，不声称可直接在所有 SDK 上编译。
+
+## 教材结构
+
+新小节标题为“全局、共享与常量内存：先判断数据该放哪里”。内容按以下顺序组织：
+
+1. **三类内存对照表**：比较作用域、读写方式、典型用途和典型误区。
+2. **短示例**：
+   - `global_add`：每个线程从 global 读取输入、写回输出，强调相邻线程地址。
+   - `block_sum`：先装载到 `__shared__`，再以 `__syncthreads()` 完成一次 block 内协作，强调它不能跨 block。
+   - `scale`：从 `__constant__` 读取一个小型只读参数，强调相同地址读取的广播条件。
+3. **选型流程**：先判断是否跨 block/跨 kernel 保存，再判断是否小型只读且访问相同地址，最后判断是否由同 block 多次复用；其余默认 global。
+4. **读 kernel 检查清单**：明确指针来源、谁读写、是否需要同步、相邻线程地址和边界保护。
+5. **通向课程**：global 指向 Week 4，shared/constant 指向 Week 5，并说明 Week 3 是 shared 同步的前置。
+
+## 示例约束
+
+每段示例少于 15 行，只展示一种核心行为：
+
+- `global_add` 必须保留 `if (i < n)`，避免把边界保护误教为可选项。
+- `block_sum` 必须让所有 block 内线程一致经过 barrier，避免在条件分支中使用 barrier；示例只产出 block partial，明确最终合并需要额外步骤。
+- `scale` 的 constant 只存单个标量，以避免读者把“常量数组索引”与“广播访问”混为一谈。
+
+## Quiz
+
+新增四道选择题，覆盖最常见的错误判断：
+
+1. `__syncthreads()` 只同步同一个 block，不能协调不同 block。
+2. global memory 中相邻线程访问相邻元素更有利于合并访问；正确性不等于高带宽。
+3. shared memory 的价值来自同一个 block 对同一份数据的复用，而不是自动缓存全部 global 数据。
+4. constant memory 适合小型只读且同一批线程经常读取相同地址的参数；每线程不同地址的访问不应假定有广播收益。
+
+所有题目补齐 `id`、`deck`、`type`、`question`、`choices`、`answerIndex`、`answer`、`pitfall` 和 `source`。题目归入 Week 4/5 关联的 deck，避免把课程阶段混淆。
+
+## 验收
+
+1. 新内容不与 Week 4/5 讲义重复大段实验说明，且明确链接到两周的学习材料。
+2. 代码块、表格和流程在 Markdown 中可读，示例不含设备无关的性能数字。
+3. `node scripts/check-docs.js` 通过，确认题目结构、来源路径和本地链接有效。
+4. `git diff --check` 通过。

scripts/check-docs.test.js

@@ -31,7 +31,7 @@ const commentLink = (href) => (html) => html.replace(/<a\b[^>]*>[\s\S]*?<\/a>/g,
   (link) => link.includes(`href="${href}"`) ? `<!--${link}-->` : link);
 
 test("baseline validates all pages and 136 local links", () => {
-  assert.deepEqual(runChecker(), ["quiz questions: 150", "learning materials: ok", "knowledge pages: 9", "local links: 136"]);
+  assert.deepEqual(runChecker(), ["quiz questions: 154", "learning materials: ok", "knowledge pages: 9", "local links: 136"]);
 });
 
 const invalidCases = [