GPU / CUDA 基础（AI Infra 方向面试对话级掌握）

目标：能就 GPU 编程聊一个来回，不是考到就哑。重点：为什么 GPU 快、内存模型、线程模型、核函数、访存与算力墙、常用优化。

一、为什么 GPU 快（解题框架）

GPU = 吞吐优先、大量简单核：几千个小核同时跑，用并行掩盖延迟。CPU few fat cores（低延迟），GPU many thin cores（高吞吐）。适合数据并行（同样操作跑海量数据：矩阵、卷积）。

二、编程模型（CUDA 词汇表，会用一词即显懂）

• Kernel：__global__ 函数，在 GPU 执行；调用 kernel<<<grid, block>>>(args)。
• 线程组织：grid（所有线程）→ block（线程块）→ thread。
• 内置索引：blockIdx, threadIdx, blockDim, gridDim。线程怎么对应数据就是手写部分。
• 内存层级：
  • 全局内存 global（大、慢、DRAM）
  • 共享内存 shared（shared，块内共享，快，需手动同步 __syncthreads）
  • 寄存器（最快）
  • 常量/纹理
• TPC/SM：SM 是执行单元，block 调度到 SM；一个 SM 里分 warp（32 线程一次执行）。 一句话给面试官：「block 内线程共享 shared memory、需 barrier 同步；warp 执行的发散会串行化不同路径。」

三、访存规则（写 kernel 必不踩）

1. 合并访问（coalescing）：相邻线程应访问相邻地址（连续数组），一次事务读满一行。违反则带宽炸（每线程单独取）。→ 数据布局几乎 SoA，索引 a[threadIdx.x] 而非跳步。
2. 对行主矩阵，c[row*w+col] 让 col=threadIdx 连续即可。
3. 需要跨线程复用 → 先拷到 shared memory（手动 tile 复用，等价 cache 但软件管理），配合 __syncthreads。
4. 别让线程 if 发散太随机（branch divergence 浪费 warp）。

四、CPU-GPU 与内存传输

• cudaMemcpy Host↔Device 走 PCIe/NVLink，慢 → 别每步都传，尽量 传 1 次、算、再取。
• cudaMalloc/cudaFree 慢（常复用 buffer / 用 stream 池）。
• 异步/stream 可重叠 copy 与 compute。

五、黄金优化顺序（背给面试官）

1. 先看 kernel 是 内存带宽受限 还是 计算受限（profiler：Nsight Compute/VTune/rocprof；成熟机型 memory bound 极常见）。带宽受限 → 减访存（合并访问、tile 复用、half/fp8 精度）。
2. 计算受限 → 提高 occupancy、减少寄存器、向量化 float4 一次取 4。
3. 调 grid/block 大小让 SM 满载、用 shared memory 复用、减少同步/分支。
4. 用 warp 聚合减少 bank conflict。 （讲“用 profiler 数据判断 CPU 段 vs 数据传输 vs kernel 谁占大头”最显专业。）

六、常见问答要义

• “GPU 比 CPU 快吗？”→ 取决于并行度与访存强度；标量小任务是 CPU。答案强调“吞吐 vs 延迟 + 数据规模 + 传输开销”。
• “为什么开你加的核数 CPU 吞吐不动但 GPU 起飞？”→ IO/带宽墙/一致性流量 vs 大规模数据并行掩盖。
• NCCL：多卡通信库（见 hpc-network）。

七、一句自我定位

「我做宿主侧 + Nsight/vtune 分析，能写 elementwise/reduction/tile kernel，清楚带宽墙与合并访问、shared 复用和 block 同步，知道何时是带宽受限还是算力受限。」