CUDA 把并行线程组织成三层:
grid
├── block
│ ├── thread
│ ├── thread
│ └── ...
├── block
│ └── ...
└── ...- thread —— 最小执行单位,跑一份 kernel 代码
- block —— 一组 thread,能共享 shared memory + 互相 __syncthreads
- grid —— 一组 block,block 之间不能直接通信(只能通过 global memory)
1. 启动一个 kernel
my_kernel<<<gridDim, blockDim>>>(args...);gridDim 和 blockDim 都是 dim3(可以是 1D / 2D / 3D)。
my_kernel<<<10, 256>>>(...); // 1D:grid=10 个 block,每 block 256 thread
my_kernel<<<dim3(4, 4), dim3(16, 16)>>>(...); // 2D:grid=4×4=16 个 block,每 block 16×16=256 thread
my_kernel<<<dim3(2, 2, 2), 64>>>(...); // 3D grid,1D block总线程数 = gridDim.x × gridDim.y × gridDim.z × blockDim.x × blockDim.y × blockDim.z。
2. 在 kernel 里查询自己的位置
__global__ void my_kernel() {
// block 在 grid 里的坐标
int bx = blockIdx.x;
int by = blockIdx.y;
// thread 在 block 里的坐标
int tx = threadIdx.x;
int ty = threadIdx.y;
// block / grid 的尺寸(整个 kernel 一致)
int bw = blockDim.x;
int gw = gridDim.x;
}5 个内建变量:blockIdx、threadIdx、blockDim、gridDim、warpSize。前四个都是 dim3,带 .x .y .z;warpSize 是常量 32(从 Tesla 一直到 Blackwell 都是 32)。
3. 三层各自的内存可见性
| 内存类型 | 谁能访问 | 速度 |
|---|---|---|
| 寄存器 | 单 thread | 最快 |
局部内存(.local) | 单 thread | 慢(DRAM) |
共享内存(__shared__) | 一个 block 内所有 thread | 快(SRAM) |
全局内存(cudaMalloc) | 所有 block,所有 grid | 慢 |
常量内存(__constant__) | 所有 thread,只读 | 缓存优化 |
关键认知:block 内可以同步、可以共享 shared memory;block 之间不行。这就是为什么很多算法要”按 block 拆分”,每个 block 完成一小块独立的工作。
4. 块内同步:__syncthreads()
__global__ void example() {
__shared__ float buf[256];
int tid = threadIdx.x;
buf[tid] = global_in[tid]; // 每 thread 写一个元素
__syncthreads(); // ★ 等整个 block 都写完
// 现在可以读别 thread 写的数据了
float v = buf[(tid + 1) % 256];
}__syncthreads() 是 block 内屏障——所有 thread 都到达这一行才一起继续。跨 block 没有对应物——想跨 block 同步只能拆成两个 kernel 启动。
5. 为什么不让 block 之间通信
设计原则:block 必须能独立调度、独立完成。
GPU 一次能在所有 SM 上同时跑几百个 block,调度顺序未定义。block A 可能先跑、block B 可能后跑,也可能同时跑。如果你假设 “block A 先写,block B 后读”,会出错——它们随时可能交换顺序。
这个限制看起来不便,但带来的好处是 kernel 自动可扩展:同一段代码在 GTX 1050(只有几个 SM)上能跑,在 H100(132 个 SM)上也能跑,只是后者并行度更高、跑得更快。
6. 一个完整例子:向量加法
__global__ void vec_add(const float* a, const float* b, float* c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= n) return; // tail guard,见下一篇
c[tid] = a[tid] + b[tid];
}
int main() {
int n = 1'000'000;
int block = 256;
int grid = (n + block - 1) / block; // 向上取整
vec_add<<<grid, block>>>(a_d, b_d, c_d, n);
cudaDeviceSynchronize();
}n = 1_000_000,block = 256 → grid = 3907。一共启动 3907 × 256 = 1,000,192 个线程,处理 100 万元素 —— 多出来的 192 个线程被 tail guard 拦住直接 return。
7. 一句话
thread / block / grid 是 CUDA 的三层并行抽象。block 内能同步 + 共享 shared memory,block 间只能通过 global memory 通信。这个限制让 kernel 自动可扩展到任意大小的 GPU。
下一篇:warp ——硬件真正调度的单位(不是 block,也不是 thread)。