<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 1 · 起步02 / 15约 22 分钟cuda/02_vector_add.cu

线程层级：grid / block / thread 的索引算术

把一维、二维数据映射到线程上，以及为什么每个 kernel 都要写边界检查

学完这一课你会
  • 熟练推导全局线程索引，包括一维和二维情形
  • 知道为什么 grid 尺寸要向上取整，以及随之而来的越界风险
  • 掌握 grid-stride loop 这个比裸索引更健壮的写法
  • 建立「线程 ↔ 数据元素」的映射直觉

CUDA 的并行模型是一个两级层次结构：grid 由若干 block 组成，block 由若干 thread 组成。这个划分不是为了好看——它直接对应硬件：一个 block 会被整体分配到一个 SM 上执行，block 内的线程可以共享 shared memory、可以用 __syncthreads() 互相等待；而不同 block 之间什么都共享不了。

                        grid
  ┌───────────────┬───────────────┬───────────────┐
  │   block 0     │   block 1     │   block 2     │
  │ t0 t1 t2 t3   │ t0 t1 t2 t3   │ t0 t1 t2 t3   │
  └───────────────┴───────────────┴───────────────┘
    0  1  2  3      4  5  6  7      8  9 10 11     ← 全局索引

  global = blockIdx.x * blockDim.x + threadIdx.x
             ↑             ↑              ↑
        我是第几个块   每块多少线程    块内第几个线程
grid(3) × block(4) 的线程编号
▸打开索引可视化实验室与其死记公式，不如亲手拖动 gridDim / blockDim，点任意一个线程看它对应哪个数据元素。二维映射尤其建议在这里点几下再往下看。

一维：向量加法

最经典的 kernel
1__global__ void vecAdd(const float* a, const float* b, float* c, int n) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3    if (i < n) {              // 边界检查：不能省4        c[i] = a[i] + b[i];5    }6}

为什么一定要写 if (i < n)？因为 block 的线程数通常取 128/256 这类 2 的幂，而 n 是任意值。要覆盖所有元素，grid 尺寸必须向上取整，于是最后一个 block 里必然有一部分线程的索引超出了数组范围。少了这行检查，就是一次显存越界写——轻则结果错乱，重则踩坏别的数据结构。

向上取整的 grid 尺寸
1int n = 1'000'000;2int threads = 256;3int blocks = (n + threads - 1) / threads;   // = ceil(n / threads) = 39074 5vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);6// 3907 * 256 = 1'000'192 个线程，最后 192 个被 if 挡住

grid-stride loop：更健壮的写法

「一个线程处理一个元素」在数据量特别大或特别小的时候都不理想。grid-stride loop 让 grid 尺寸和数据规模解耦：不管你启动多少线程，每个线程都用固定步长把整个数组扫一遍。这样同一份 kernel 可以用一个针对硬件调优过的 grid 尺寸跑任意规模的数据。

grid-stride loop
1__global__ void vecAddStride(const float* a, const float* b, float* c, int n) {2    int stride = gridDim.x * blockDim.x;         // 整个 grid 的线程总数3    for (int i = blockIdx.x * blockDim.x + threadIdx.x;4         i < n;5         i += stride) {6        c[i] = a[i] + b[i];7    }8}
  • ▸规模无关：n 是 1000 还是 10 亿，同一个 grid 配置都能正确跑完。
  • ▸访问依然合并：相邻线程在每一轮迭代里访问的仍是相邻地址，不破坏内存合并（下一阶段会讲）。
  • ▸便于调试：把 grid 设成 <<<1, 1>>> 就退化成串行执行，方便你对拍结果。
  • ▸复用寄存器：循环外的常量计算只做一次，摊薄到多个元素上。

二维：图像和矩阵

处理矩阵时用二维 grid 更自然。注意一个约定：threadIdx.x 应该对应行内连续的方向（也就是列号），因为 warp 内相邻的线程是按 .x 优先编号的，只有让 .x 走连续内存，访问才能合并。这一点写反了，性能可能差 5 倍以上。

二维索引：注意 x 对应列
1__global__ void scaleMatrix(float* m, int rows, int cols, float k) {2    // x → 列（内存连续方向），y → 行3    int col = blockIdx.x * blockDim.x + threadIdx.x;4    int row = blockIdx.y * blockDim.y + threadIdx.y;5 6    if (row < rows && col < cols) {7        m[row * cols + col] *= k;     // 行主序：同一 warp 覆盖连续的 col8    }9}10 11// host 侧12dim3 block(32, 8);                                    // 256 线程，x 维正好一个 warp13dim3 grid((cols + block.x - 1) / block.x,14          (rows + block.y - 1) / block.y);15scaleMatrix<<<grid, block>>>(d_m, rows, cols, 2.0f);

block 大小怎么选

block 尺寸评价
32太小。每个 SM 能驻留的 block 数有上限（通常 16~32），线程总量上不去，占用率被卡死
128 / 256默认首选。绝大多数 kernel 从这里起步，兼顾占用率与调度灵活性
512适合每线程寄存器少、需要大块 shared memory 协同的 kernel
1024硬件上限。寄存器压力极大，通常反而更慢，除非算法确实需要
非 32 倍数永远不要。比如 100，会让每个 block 的最后一个 warp 有 28 个线程空转

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。