执行层次与索引

Grid（一维/二维/三维）
└─ Block（一维/二维/三维）
   └─ Thread

一维全局索引：i = blockIdx.x * blockDim.x + threadIdx.x
二维全局索引：x = blockIdx.x * blockDim.x + threadIdx.x
              y = blockIdx.y * blockDim.y + threadIdx.y

线程坐标负责“谁处理哪个数据”，数据布局负责“这个数据在内存哪里”。二者应一起设计，让同一 warp 的相邻 lane 尽量访问相邻地址。

一维、二维与三维映射

一维：数组与元素级算子

__global__ void saxpy(int n, float a, const float* x, float* y) {
  for (int i = blockIdx.x * blockDim.x + threadIdx.x;
       i < n;
       i += blockDim.x * gridDim.x) {
    y[i] = a * x[i] + y[i];
  }
}

这是 grid-stride loop。优点：

• grid 不必大到为每个元素创建一个逻辑线程；
• 能处理大于 grid 维度的输入；
• 便于控制 block 数和复用线程；
• 每个 warp 在每轮仍访问连续元素。

初始 block size 可从 128 或 256 开始实测。网格可取足够多的 block 覆盖所有元素，或取 若干倍 SM 数 再依赖 grid-stride loop；后者适合控制启动规模，但倍率必须测量。

二维：图像、矩阵与网格

__global__ void rgb_to_gray(const uchar3* src, unsigned char* dst,
                            int width, int height) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if (x >= width || y >= height) return;

  uchar3 p = src[y * width + x];
  dst[y * width + x] = static_cast<unsigned char>(
      0.299f * p.x + 0.587f * p.y + 0.114f * p.z);
}

dim3 block(32, 8);
dim3 grid((width + block.x - 1) / block.x,
          (height + block.y - 1) / block.y);

x 放在连续维，使相邻 lane 横向移动。32×8 与 16×16 都是 256 线程，但前者往往更自然地对齐行主序访存；最终仍以实际访问模式和指标为准。

三维：体素与 PDE 网格

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;
size_t i = (static_cast<size_t>(z) * ny + y) * nx + x;

乘法前转换成 size_t，避免大数组的 32 位整数溢出。三维 block 不必在每维都很大，总线程数仍受设备的 maxThreadsPerBlock 限制。

Block 大小的资源权衡

block threads 增加
  ├─ 每 block warp 增多，block 内协作范围更大
  ├─ block 总寄存器用量增加
  ├─ 往往需要更多 shared memory
  └─ 每 SM 可同时驻留的 block 可能减少

常用经验只是起点：

• thread 数取 warp size 的整数倍；
• 128–256 常是合理首测范围；
• grid 中 block 数应明显多于 SM 数，除非是 persistent kernel；
• 包含 block barrier 的 kernel 通常希望每 SM 不止一个 active block；
• 用 occupancy API、编译器资源报告和 Nsight Compute 验证。

int min_grid = 0, suggested_block = 0;
cudaOccupancyMaxPotentialBlockSize(
    &min_grid, &suggested_block, saxpy, 0, 0);

Occupancy API 给的是资源约束下的候选配置，不知道你的访存、分支和算法结构，不能替代 benchmark。

Kernel 函数限定符

常见变量限定符：

__host__ __device__ 函数会分别编译 host 与 device 版本。函数体中使用仅一侧支持的代码时，需要条件编译或清晰分层；否则模板错误和行为差异会很难读。

Warp 级原语

现代 CUDA 的 shuffle/vote 操作都带 active mask：

unsigned mask = __activemask();
float v = ...;
for (int offset = warpSize / 2; offset > 0; offset >>= 1) {
  v += __shfl_down_sync(mask, v, offset);
}

不要写“warp 天然 lockstep，所以不需要同步”的旧式代码。Volta 起的独立线程调度让隐式 warp-synchronous 假设尤其危险；使用带 _sync 的原语、正确 mask 和显式同步。

Cooperative Groups

Cooperative Groups 把协作范围变成显式对象：

#include <cooperative_groups.h>
namespace cg = cooperative_groups;

__global__ void kernel(float* x) {
  cg::thread_block block = cg::this_thread_block();
  auto tile = cg::tiled_partition<32>(block);

  // block.thread_rank(), block.num_threads(), block.sync()
  // tile.thread_rank(), tile.sync(), cg::reduce(...)
}

优势不是语法更长，而是被调用函数能声明“所有组成员必须参与”的合同，避免只让半个 block 调用内部 barrier 的死锁：

__device__ float reduce_group(cg::thread_block g, float x) {
  // 调用方必须传入并共同参与同一个 group。
  ...
  g.sync();
  ...
}

不规则并行策略

首先改变任务到 warp/block 的映射，而不是立即引入复杂的工作窃取。分桶（按工作量或数据特征分类）经常能同时改善负载均衡与分支一致性。

边界线程与提前返回

这段是安全的：

int i = ...;
if (i >= n) return;
out[i] = f(in[i]);

但如果后面有 block barrier，提前返回会导致并非所有 block 线程抵达同步点：

// 错误模式
if (i >= n) return;
shared[threadIdx.x] = in[i];
__syncthreads();

应让所有线程参与 barrier，只对加载和写回做谓词保护：

shared[threadIdx.x] = (i < n) ? in[i] : 0.0f;
__syncthreads();
if (i < n) out[i] = use(shared);

设计一个 Kernel 的检查表

1. 并行单元是什么：元素、行、tile、边、样本还是批次？
2. 一个 thread/warp/block 各负责多少工作？
3. 同一 warp 的访问是否连续、对齐、尽量少事务？
4. 数据是否在线程、warp、block 内复用？放在哪一级存储？
5. 是否所有需要参与 barrier 的线程都会到达？
6. grid 是否足够大？尾部 block 是否正确？
7. 寄存器和 shared memory 会限制多少 active blocks？
8. 有 CPU reference、边界输入和随机输入测试吗？

官方参考

• Writing SIMT Kernels
• Cooperative Groups
• CUDA C++ Best Practices: Execution Configuration