Guadazi
0%

CUDA Kernel 代码阅读能力

Posted on Edited on In

CUDA Kernel 代码阅读能力

一句话定位：不要求会写复杂 Kernel，但要能读懂三类最基础的代码——索引计算、边界判断、访存模式，这三点决定了一段 Kernel 是否正确、是否高效。

1. 全局索引计算：定位”我是谁”

每个线程要先知道自己该处理数据的哪一部分，标准写法：

1
int idx = blockIdx.x * blockDim.x + threadIdx.x;
  • blockIdx.x：当前线程所在的 Block 编号；blockDim.x：每个 Block 里的线程数；threadIdx.x：线程在 Block 内的局部编号。
  • 这行代码把”块内局部编号”转换成”全局唯一编号”，与大数据里”分区编号 × 分区大小 + 分区内偏移量 = 全局偏移”的思路完全一致。

2. 边界判断：防止越界访问

线程总数通常是 Block/Thread 配置向上取整的结果，会超过实际数据长度，因此几乎每个 Kernel 都要有：

1
2
3
if (idx < n) {
    // 实际计算逻辑
}

漏掉这行是最常见的越界访问 Bug 来源（读到脏数据或直接 crash）。

3. 三段典型代码要读懂

  • 向量加法（最简单）：每个线程处理一个元素，c[idx] = a[idx] + b[idx]，重点看索引计算与边界判断。
  • 朴素矩阵乘：每个线程计算结果矩阵的一个元素，需要一个内层循环遍历公共维度；重点看二维索引 row = blockIdx.y*blockDim.y+threadIdx.y、col = blockIdx.x*blockDim.x+threadIdx.x 的推导，以及这种朴素写法的访存是否高效。
  • Reduction（归约，如求和）：多个线程协作把一组数据归约成一个值，典型做法是先在共享内存里做树形归约（每轮线程数减半），再跨 Block 归约；重点看共享内存的读写同步（__syncthreads()）与树形归约的下标规律。

4. 访存是否合并（Coalesced Access）

  • 合并访问：同一个 Warp（32 个线程）里的线程访问连续的显存地址，硬件可以合并成一次内存事务，效率最高。
  • 非合并访问：线程访问地址跳跃或不连续（如按列访问按行存储的矩阵），会拆成多次内存事务，显存带宽利用率骤降。
  • 判断方法：看索引表达式里线程编号（threadIdx.x）是否直接映射到内存地址的最低维、连续位置。

5. 面试落点

  • 产出物：能口头讲清一段给定 CUDA Kernel 的执行逻辑（每个线程做什么）与访存模式（是否合并访问）。
  • 可迁移话术：矩阵乘的朴素实现（未做 Tiling/共享内存复用）对应大数据里”没有做本地聚合就直接 Shuffle”的低效模式；Reduction 的树形归约思路与 Spark 的 treeReduce（相对于普通 reduce 减少 Driver 端压力）是同一个思想的不同实现层。

参考：《CUDA C++ Programming Guide》；NVIDIA CUDA Samples 官方仓库