目标
GPU 上实现 O(n) 粒子模拟
一句话总结：同一个粒子模拟问题（HW2 的 OpenMP / MPI 版本），这次在 GPU 上并行化——你只能修改 gpu.cu 中的 simulate_one_step 函数，在 NVIDIA A100 GPU 上跑出 O(n) 的性能。
这是你第一次用 CUDA 写"真正"的并行程序。与 CPU 并行不同，GPU 编程需要你同时思考线程布局、内存访问模式和同步策略。
| 任务 | 对应课程 | 核心概念 |
|---|---|---|
| 理解 GPU 硬件模型 | Lecture 9 | SM、Warp、SIMT、延迟隐藏 |
| 设计线程布局 | Lecture 9 | Grid / Block / Thread 映射 |
| 实现 simulate_one_step | Lecture 10 | 数据并行 kernel |
| 优化内存访问 | Lecture 10 | 合并访问、shared memory |
| 性能评测与对比 | Lecture 10 | 与 OpenMP / MPI 版本对比 |
| 撰写报告 | — | log-log 图、扩展性分析 |
__global__、__device__、__host__ 修饰符的含义<<<blocks, threads>>> 执行配置语法threadIdx、blockIdx、blockDim、gridDim 的含义cudaMalloc / cudaMemcpy / cudaFree 的用法__syncthreads() 的语义和限制回顾 HW2 的粒子模拟：每个时间步，每个粒子根据与其他粒子的相互作用更新位置和速度。核心循环是每个粒子遍历所有其他粒子——O(n²) 的朴素算法。理解这一步是设计 GPU kernel 的前提。
最自然的映射是一个线程处理一个粒子：
// 全局线程 ID 计算
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < n) {
// 粒子 i 的更新逻辑
}选择合适的 block size（通常 128 或 256 线程/block），计算所需的 block 数量。
你需要：
cudaMalloc 在 device 上分配对应空间cudaMemcpy 将数据拷贝到 GPUGPU 性能的关键是内存合并（coalesced access）：
用不同粒子数（100K - 10M）测试，绘制 log-log 图，与 HW2 的 OpenMP / MPI 版本对比。注意：小粒子数看不出差异，必须用大规模数据。
忘记边界检查： 当 n 不是 blockDim 的整数倍时，最后一个 block 会有多余线程。必须用 if (i < n) 保护，否则会访问越界。
Host/Device 指针混用： 不能在 host 代码中解引用 device 指针，反之亦然。这是最常见的运行时错误。
分支发散： 如果 warp 内线程走不同分支（如某些粒子需要特殊处理），性能会大幅下降。尽量让 warp 内线程执行相同路径。
非合并访问： 如果线程 i 访问 data[i * stride] 且 stride > 1，会导致非合并访问，带宽利用率可能只有 1/32。
同步位置错误： __syncthreads() 必须在所有线程都能到达的位置调用。放在条件分支内可能导致死锁。
| 概念 | 含义 | HW2-3 中的对应 |
|---|---|---|
| Thread | 单个执行单元 | 处理一个粒子的更新 |
| Warp | 32 个线程的调度组 | 锁步执行，注意分支发散 |
| Block | 驻留在一个 SM 上的线程组 | 共享 shared memory，可同步 |
| Grid | 所有 block 的集合 | 一次 kernel 启动的全部工作 |
| 内存 | 延迟 | 用途 |
|---|---|---|
| 寄存器 | ~1 cycle | 线程私有变量（循环计数器、累加器） |
| Shared Memory | ~5 cycles | Block 内线程共享数据（tile 缓存） |
| Global Memory | ~400-800 cycles | 粒子位置/速度的主存储 |
HW2-3 是 GPU 编程的入门体验。后续课程中你会见到更复杂的 GPU 应用：
卡住时的正确顺序： 复习 Lecture 9-10 → 读 handout → 从最小 kernel 开始 → 用 printf 调试 → 用 nsight compute 分析。
问题 1：GPU 中 warp 的大小是多少？当 warp 内线程走不同分支时会发生什么？
问题 2：为什么 GPU 编程强调"合并内存访问"（coalesced access）？
问题 3：__syncthreads() 的同步范围是什么？
问题 4：为什么 HW2-3 要求用 100K-10M 的大粒子数？
GPU 上实现 O(n) 粒子模拟
线程布局 + 合并访问 + 同步
与 OpenMP / MPI 版本性能对比
log-log 图 + 设计解释 + 分析