Docs/learning roadmap and cuda reference

Agent.md

@@ -22,7 +22,8 @@
 3. Week 主线练习放 `code/weekN/exercises.md`；刷题索引放 `code/weekN/question-bank/`；快速复习材料可放 `code/weekN/concept-review.md`。
 4. 性能数字不要伪造。需要实测的表格可以留 `?`，并说明记录到对应 `notes/weekN.md`。
 5. MUSA 与 CUDA 很像，但不要默认性能经验完全相同。尤其注意 MUSA warp size 通常按 128 讨论。
-6. 新增知识点、示例、case、概念文档或复习材料时，必须同步评估是否需要给 `docs/index.html` 增加对应测试题；如果不增加题目，需要在提交或 PR 描述里说明原因。
+6. 新增知识点、示例、case、概念文档或复习材料时，必须在同一次修改中给 `docs/index.html` 的 MUSA Quiz 增加对应测试题；不得只更新学习材料而把题库同步留到以后。基础概念（如 thread、block、grid、warp、索引计算、内存层级、同步、访存）、代码写法、易错点和面试问法也必须纳入题库。题目必须覆盖新增内容的关键概念或易错点，并补齐 `answer`、`pitfall`、`source` 字段。
+7. 每次新增或修改学习材料后，必须运行题库和文档校验；如果当前环境没有 MUSA SDK，至少运行静态校验，并在结果中明确不能声称 MUSA 编译通过。
 
 ## 构建与运行
 
@@ -80,6 +81,7 @@ Mac 本地没有 `mcc` 时，可以本地编辑、远端运行：
 ## Git 与验证
 
 - 提交前运行 `git diff --check`。
+- 每次新增学习内容后运行 `node scripts/check-docs.js`，确认 Quiz 题目结构、题目来源路径和学习材料完整。
 - 文档改动应检查 Markdown 链接、图片路径和代码块是否闭合。
 - MUSA 示例在没有 MUSA SDK 的 Mac 上不能声称“编译通过”；应明确说明只做了静态检查或远程验证。
 - 不要把用户已有的未提交修改混入无关 PR；提交前按目标路径检查 staged diff。

code/week1/06_async_kernel.mu

@@ -16,6 +16,7 @@
 //
 //      kernel<<<g, b>>>(args);
 //      // ↑ 这一句几乎不耗时（只是把"启动请求"丢进队列）
+//      // 异步执行
 //      // CPU 不会等 GPU 算完，立刻继续往下跑
 //
 //  好处：CPU 可以在 GPU 计算时同时做别的事（比如准备下一批数据），
@@ -32,6 +33,7 @@
 // └──────────────────────────────────────────────────────────────────────────
 //
 //      kernel<<<>>>(d_out, ...);   // GPU 还在算
+//      // 下面这个是在 Host 上操作，所以不能安全读取 Device 上的数据
 //      printf("%f\n", h_out[0]);   // ← 读的是脏数据，因为根本没拷回
 //
 //  正确写法：

code/week1/README.md

@@ -84,3 +84,7 @@ Grid / block / thread 的层级、索引换算和边界判断容易混淆，先
 LeetGPU Easy 题，按必做、进阶和选做分层。
 
 CUDA_Freshman 对照案例和迁移计划见 [`../../docs/cuda-example-map.md`](../../docs/cuda-example-map.md)。
+
+## CUDA reference（精选）
+
+来自 *Professional CUDA C Programming* 的基础 `.cu` 示例已放在 [`cuda-reference/`](cuda-reference/)。它们保留 CUDA API，使用目录内 `make BACKEND=cuda|musa` 区分编译后端；MUSA 仅记录兼容路径，未在当前环境声称实编通过。逐文件 target、限制和运行命令见 [`cuda-reference/README.md`](cuda-reference/README.md)。

code/week1/cuda-reference/Makefile

@@ -0,0 +1,58 @@
+# Common build contract for the CUDA reference examples.
+# Source targets are discovered under chapterXX/ as the snapshot is imported.
+
+BACKEND ?= cuda
+NVCC ?= nvcc
+MCC ?= mcc
+MUSA_HOME ?= /usr/local/musa
+CUDA_ARCH ?=
+MUSA_ARCH ?=
+TARGET ?= all
+
+# Keep this list explicit until reference sources are imported.
+SOURCES := $(shell find . -type f -name '*.cu' -not -path './build/*' | sort)
+target_name = $(subst /,__,$(patsubst ./%,%,$(patsubst %.cu,%,$(1))))
+TARGETS := $(foreach source,$(SOURCES),$(call target_name,$(source)))
+
+ifeq ($(BACKEND),cuda)
+  COMPILER := $(NVCC)
+  BACKEND_FLAGS := -O2 -lineinfo $(if $(CUDA_ARCH),-arch=$(CUDA_ARCH),)
+  BACKEND_DEFINES :=
+else ifeq ($(BACKEND),musa)
+  COMPILER := $(MCC)
+  BACKEND_FLAGS := -O2 -lineinfo --musa-path=$(MUSA_HOME) -mtgpu $(if $(MUSA_ARCH),--offload-arch=$(MUSA_ARCH),)
+  BACKEND_DEFINES := -DCUDA_REFERENCE_MUSA
+else
+  $(error Unknown BACKEND '$(BACKEND)'; choose cuda or musa)
+endif
+
+ifneq ($(TARGET),all)
+  ifeq ($(filter $(TARGET),$(TARGETS)),)
+    $(error Unknown TARGET '$(TARGET)'; available targets: all$(if $(TARGETS), $(TARGETS), none yet))
+  endif
+endif
+
+.DEFAULT_GOAL := $(if $(filter all,$(TARGET)),all,$(TARGET))
+.PHONY: all clean optional
+
+define make-target
+$(call target_name,$(1)): build/$(call target_name,$(1))
+
+build/$(call target_name,$(1)): $(1)
+	@mkdir -p build
+	$(COMPILER) $(BACKEND_FLAGS) $(BACKEND_DEFINES) -Icommon $$< -o $$@
+endef
+$(foreach source,$(SOURCES),$(eval $(call make-target,$(source))))
+
+ifeq ($(strip $(TARGETS)),)
+all:
+	@echo "$(COMPILER) $(BACKEND_FLAGS) <future-target>.cu -o <future-target>"
+	@echo "No reference source targets imported yet; this is a dry-run build skeleton."
+else
+all: $(TARGETS)
+endif
+
+optional:
+
+clean:
+	@rm -rf ./build

code/week1/cuda-reference/README.md

@@ -0,0 +1,27 @@
+# Week 1 CUDA reference：基础执行模型
+
+来源：[`kriegalex/wrox-pro-cuda-c`](https://github.com/kriegalex/wrox-pro-cuda-c)，固定 commit `63825d64683b644198dd9cb0d4d472d6914d4f72`，MIT license。源码保留 `.cu` 和 CUDA API；本目录的 Makefile 用 `BACKEND=cuda` 选择 `nvcc`，用 `BACKEND=musa` 选择 `mcc -mtgpu`。MUSA 状态均为“未在本机实编”，兼容性取决于 SDK、MUSA Mapping、架构和 API 支持。
+
+| 文件 | 主题 / 来源 | target | CUDA 状态 | MUSA 状态 / 限制 |
+|---|---|---|---|---|
+| `chapter01/hello.cu` | kernel 与 GPU printf / Ch1 | `chapter01__hello` | 预计可用 `nvcc`；未实编 | 基础语法通常可映射；未实编 |
+| `chapter02/checkDeviceInfor.cu` | 设备属性 / Ch2 | `chapter02__checkDeviceInfor` | 预计可用 `nvcc`；未实编 | 属性字段随 runtime 变化；未实编 |
+| `chapter02/checkDimension.cu` | grid/block/thread 维度 / Ch2 | `chapter02__checkDimension` | 预计可用 `nvcc`；未实编 | 基础 API；未实编 |
+| `chapter02/checkThreadIndex.cu` | 线程索引 / Ch2 | `chapter02__checkThreadIndex` | 预计可用 `nvcc`；未实编 | 基础 API；未实编 |
+| `chapter02/defineGridBlock.cu` | launch 配置 / Ch2 | `chapter02__defineGridBlock` | 预计可用 `nvcc`；未实编 | 受设备最大 block 限制；未实编 |
+| `chapter02/sumArraysOnGPU-small-case.cu` | 单 block 向量加 / Ch2 | `chapter02__sumArraysOnGPU-small-case` | 预计可用 `nvcc`；未实编 | 基础 kernel；未实编 |
+| `chapter02/sumArraysOnGPU-timer.cu` | CPU/GPU 计时 / Ch2 | `chapter02__sumArraysOnGPU-timer` | 预计可用 `nvcc`；未实编 | 计时结果不能跨设备外推；未实编 |
+| `chapter02/sumMatrixOnGPU-1D-grid-1D-block.cu` | 1D grid/block 矩阵加 / Ch2 | `chapter02__sumMatrixOnGPU-1D-grid-1D-block` | 预计可用 `nvcc`；未实编 | 观察 MUSA warp/block 差异；未实编 |
+| `chapter02/sumMatrixOnGPU-2D-grid-1D-block.cu` | 2D grid/1D block / Ch2 | `chapter02__sumMatrixOnGPU-2D-grid-1D-block` | 预计可用 `nvcc`；未实编 | 设备限制依赖 SDK；未实编 |
+| `chapter02/sumMatrixOnGPU-2D-grid-2D-block.cu` | 2D grid/2D block / Ch2 | `chapter02__sumMatrixOnGPU-2D-grid-2D-block` | 预计可用 `nvcc`；未实编 | 设备限制依赖 SDK；未实编 |
+
+## 运行
+
+```bash
+cd code/week1/cuda-reference
+make BACKEND=cuda TARGET=chapter01__hello
+./build/chapter01__hello
+make BACKEND=musa MUSA_ARCH=mp_31 TARGET=chapter01__hello
+```
+
+`make` 默认会尝试编译本目录全部 `.cu`；没有 CUDA/MUSA SDK 时只运行 `make -n` 做命令检查。公共头只有 `code/week1/cuda-reference/common/common.h` 一份，源码通过 `-Icommon` 使用。

code/week1/cuda-reference/chapter01/hello.cu

@@ -0,0 +1,21 @@
+#include "../common/common.h"
+#include <stdio.h>
+
+/*
+ * A simple introduction to programming in CUDA. This program prints "Hello
+ * World from GPU! from 10 CUDA threads running on the GPU.
+ */
+
+__global__ void helloFromGPU()
+{
+    printf("Hello World from GPU!\n");
+}
+
+int main(int argc, char **argv)
+{
+    printf("Hello World from CPU!\n");
+
+    helloFromGPU<<<1, 10>>>();
+    CHECK(cudaDeviceReset());
+    return 0;
+}

code/week1/cuda-reference/chapter02/checkDeviceInfor.cu

@@ -0,0 +1,91 @@
+#include "../common/common.h"
+#include <cuda_runtime.h>
+#include <stdio.h>
+
+/*
+ * Display a variety of information on the first CUDA device in this system,
+ * including driver version, runtime version, compute capability, bytes of
+ * global memory, etc.
+ */
+
+int main(int argc, char **argv)
+{
+    printf("%s Starting...\n", argv[0]);
+
+    int deviceCount = 0;
+    cudaGetDeviceCount(&deviceCount);
+
+    if (deviceCount == 0)
+    {
+        printf("There are no available device(s) that support CUDA\n");
+    }
+    else
+    {
+        printf("Detected %d CUDA Capable device(s)\n", deviceCount);
+    }
+
+    int dev = 0, driverVersion = 0, runtimeVersion = 0;
+    CHECK(cudaSetDevice(dev));
+    cudaDeviceProp deviceProp;
+    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
+    printf("Device %d: \"%s\"\n", dev, deviceProp.name);
+
+    cudaDriverGetVersion(&driverVersion);
+    cudaRuntimeGetVersion(&runtimeVersion);
+    printf("  CUDA Driver Version / Runtime Version          %d.%d / %d.%d\n",
+           driverVersion / 1000, (driverVersion % 100) / 10,
+           runtimeVersion / 1000, (runtimeVersion % 100) / 10);
+    printf("  CUDA Capability Major/Minor version number:    %d.%d\n",
+           deviceProp.major, deviceProp.minor);
+    printf("  Total amount of global memory:                 %.2f GBytes (%llu "
+           "bytes)\n", (float)deviceProp.totalGlobalMem / pow(1024.0, 3),
+           (unsigned long long)deviceProp.totalGlobalMem);
+    printf("  GPU Clock rate:                                %.0f MHz (%0.2f "
+           "GHz)\n", deviceProp.clockRate * 1e-3f,
+           deviceProp.clockRate * 1e-6f);
+    printf("  Memory Clock rate:                             %.0f Mhz\n",
+           deviceProp.memoryClockRate * 1e-3f);
+    printf("  Memory Bus Width:                              %d-bit\n",
+           deviceProp.memoryBusWidth);
+
+    if (deviceProp.l2CacheSize)
+    {
+        printf("  L2 Cache Size:                                 %d bytes\n",
+               deviceProp.l2CacheSize);
+    }
+
+    printf("  Max Texture Dimension Size (x,y,z)             1D=(%d), "
+           "2D=(%d,%d), 3D=(%d,%d,%d)\n", deviceProp.maxTexture1D,
+           deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1],
+           deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1],
+           deviceProp.maxTexture3D[2]);
+    printf("  Max Layered Texture Size (dim) x layers        1D=(%d) x %d, "
+           "2D=(%d,%d) x %d\n", deviceProp.maxTexture1DLayered[0],
+           deviceProp.maxTexture1DLayered[1], deviceProp.maxTexture2DLayered[0],
+           deviceProp.maxTexture2DLayered[1],
+           deviceProp.maxTexture2DLayered[2]);
+    printf("  Total amount of constant memory:               %lu bytes\n",
+           deviceProp.totalConstMem);
+    printf("  Total amount of shared memory per block:       %lu bytes\n",
+           deviceProp.sharedMemPerBlock);
+    printf("  Total number of registers available per block: %d\n",
+           deviceProp.regsPerBlock);
+    printf("  Warp size:                                     %d\n",
+           deviceProp.warpSize);
+    printf("  Maximum number of threads per multiprocessor:  %d\n",
+           deviceProp.maxThreadsPerMultiProcessor);
+    printf("  Maximum number of threads per block:           %d\n",
+           deviceProp.maxThreadsPerBlock);
+    printf("  Maximum sizes of each dimension of a block:    %d x %d x %d\n",
+           deviceProp.maxThreadsDim[0],
+           deviceProp.maxThreadsDim[1],
+           deviceProp.maxThreadsDim[2]);
+    printf("  Maximum sizes of each dimension of a grid:     %d x %d x %d\n",
+           deviceProp.maxGridSize[0],
+           deviceProp.maxGridSize[1],
+           deviceProp.maxGridSize[2]);
+    printf("  Maximum memory pitch:                          %lu bytes\n",
+           deviceProp.memPitch);
+
+    exit(EXIT_SUCCESS);
+}

code/week1/cuda-reference/chapter02/checkDimension.cu

@@ -0,0 +1,40 @@
+#include "../common/common.h"
+#include <cuda_runtime.h>
+#include <stdio.h>
+
+/*
+ * Display the dimensionality of a thread block and grid from the host and
+ * device.
+ */
+
+__global__ void checkIndex(void)
+{
+    printf("threadIdx:(%d, %d, %d)\n", threadIdx.x, threadIdx.y, threadIdx.z);
+    printf("blockIdx:(%d, %d, %d)\n", blockIdx.x, blockIdx.y, blockIdx.z);
+
+    printf("blockDim:(%d, %d, %d)\n", blockDim.x, blockDim.y, blockDim.z);
+    printf("gridDim:(%d, %d, %d)\n", gridDim.x, gridDim.y, gridDim.z);
+
+}
+
+int main(int argc, char **argv)
+{
+    // define total data element
+    int nElem = 6;
+
+    // define grid and block structure
+    dim3 block(3);
+    dim3 grid((nElem + block.x - 1) / block.x);
+
+    // check grid and block dimension from host side
+    printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);
+    printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
+
+    // check grid and block dimension from device side
+    checkIndex<<<grid, block>>>();
+
+    // reset device before you leave
+    CHECK(cudaDeviceReset());
+
+    return(0);
+}

code/week1/cuda-reference/chapter02/checkThreadIndex.cu

@@ -0,0 +1,96 @@
+#include "../common/common.h"
+#include <cuda_runtime.h>
+#include <stdio.h>
+
+/*
+ * This example helps to visualize the relationship between thread/block IDs and
+ * offsets into data. For each CUDA thread, this example displays the
+ * intra-block thread ID, the inter-block block ID, the global coordinate of a
+ * thread, the calculated offset into input data, and the input data at that
+ * offset.
+ */
+
+void printMatrix(int *C, const int nx, const int ny)
+{
+    int *ic = C;
+    printf("\nMatrix: (%d.%d)\n", nx, ny);
+
+    for (int iy = 0; iy < ny; iy++)
+    {
+        for (int ix = 0; ix < nx; ix++)
+        {
+            printf("%3d", ic[ix]);
+
+        }
+
+        ic += nx;
+        printf("\n");
+    }
+
+    printf("\n");
+    return;
+}
+
+__global__ void printThreadIndex(int *A, const int nx, const int ny)
+{
+    int ix = threadIdx.x + blockIdx.x * blockDim.x;
+    int iy = threadIdx.y + blockIdx.y * blockDim.y;
+    unsigned int idx = iy * nx + ix;
+
+    printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d,%d) global index"
+           " %2d ival %2d\n", threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y,
+           ix, iy, idx, A[idx]);
+}
+
+int main(int argc, char **argv)
+{
+    printf("%s Starting...\n", argv[0]);
+
+    // get device information
+    int dev = 0;
+    cudaDeviceProp deviceProp;
+    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
+    printf("Using Device %d: %s\n", dev, deviceProp.name);
+    CHECK(cudaSetDevice(dev));
+
+    // set matrix dimension
+    int nx = 8;
+    int ny = 6;
+    int nxy = nx * ny;
+    int nBytes = nxy * sizeof(float);
+
+    // malloc host memory
+    int *h_A;
+    h_A = (int *)malloc(nBytes);
+
+    // iniitialize host matrix with integer
+    for (int i = 0; i < nxy; i++)
+    {
+        h_A[i] = i;
+    }
+    printMatrix(h_A, nx, ny);
+
+    // malloc device memory
+    int *d_MatA;
+    CHECK(cudaMalloc((void **)&d_MatA, nBytes));
+
+    // transfer data from host to device
+    CHECK(cudaMemcpy(d_MatA, h_A, nBytes, cudaMemcpyHostToDevice));
+
+    // set up execution configuration
+    dim3 block(4, 2);
+    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
+
+    // invoke the kernel
+    printThreadIndex<<<grid, block>>>(d_MatA, nx, ny);
+    CHECK(cudaGetLastError());
+
+    // free host and devide memory
+    CHECK(cudaFree(d_MatA));
+    free(h_A);
+
+    // reset device
+    CHECK(cudaDeviceReset());
+
+    return (0);
+}