CUDA 编程模型：Thread、Block、Grid 与内存协作 / CUDA Threads, Blocks, Grids, and Cooperative Memory Access

📅 创建时间：2026-07-20 🏷️ 标签：#CUDA #Kernel #ThreadBlock #GPU编程 📚 前置知识：[[05-gpu-architecture]] 📚 相关知识：[[/02-systems-and-performance/02-computer-architecture-and-hardware/04-cuda-programming-model]] [[/02-systems-and-performance/02-computer-architecture-and-hardware/05-cuda-kernel-and-memory]]

1. Host 与 Device

CUDA 程序通常同时运行在两个世界：

text

CPU / Host                         GPU / Device
准备数据                           执行大量并行线程
分配显存        ───────────────→   读取显存
启动 Kernel                        完成计算
读取结果        ←───────────────   写回显存

CPU 负责控制流程，GPU 负责执行 Kernel。一次 GPU 加速是否值得，必须把数据传输和 Kernel 启动成本也算进去。

2. 第一个 Kernel

cpp

__global__ void vectorAdd(const float* a, const float* b,
                          float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

int threads = 256;
int blocks = (n + threads - 1) / threads;
vectorAdd<<<blocks, threads>>>(a, b, c, n);

三个重要索引：

threadIdx：线程在线程块中的位置
blockIdx：线程块在网格中的位置
blockDim：每个线程块的尺寸

3. Grid、Block 与 Thread

text

Grid
├─ Block 0
│  ├─ Thread 0
│  ├─ Thread 1
│  └─ ...
├─ Block 1
└─ ...

设计原则：

Thread 处理最细粒度工作，例如一个元素
Block 是能够共享内存并进行块内同步的线程组
Grid 表示一次 Kernel 启动的全部线程块

不同 Block 之间默认不能通过普通 Barrier 同步，因此算法需要把跨 Block 协作放到后续 Kernel 或专门机制中。

4. 线程块怎样映射到硬件

线程块被调度到 SM，内部线程按 Warp 分组。常见一维 Block 大小为 128、256 或 512，但没有对所有 Kernel 都最优的固定值。

选择 Block 大小时要考虑：

是否为 Warp 大小的整数倍
每线程寄存器数量
每 Block Shared Memory 用量
Kernel 的访存和计算特征
设备允许的最大线程数

Occupancy 高不等于性能一定高，它只表示活跃 Warp 相对于硬件上限的比例。

5. Shared Memory 协作

线程块可以先把数据从 Global Memory 搬到 Shared Memory，再重复使用：

cpp

__global__ void tiledKernel(const float* input, float* output) {
    __shared__ float tile[256];
    int global = blockIdx.x * blockDim.x + threadIdx.x;

    tile[threadIdx.x] = input[global];
    __syncthreads();

    output[global] = computeWithNeighbors(tile, threadIdx.x);
}

__syncthreads() 要求 Block 中所有仍然参与执行的线程到达同步点。把它放进线程条件不一致的分支可能导致错误或死锁。

6. 异步执行与 Stream

Kernel 启动通常对 CPU 是异步的。CUDA Stream 表示设备上的有序任务队列：

text

Stream 0：拷贝 A → Kernel A → 回传 A
Stream 1：       拷贝 B → Kernel B → 回传 B

在硬件支持、数据独立且使用异步 API 的条件下，可以实现：

数据传输与计算重叠
多个 Kernel 并发
双缓冲流水线

Pinned Host Memory 有利于异步 DMA 传输，但锁页内存是有限资源，不应无限申请。

7. 错误检查

Kernel 错误可能在后续同步时才暴露。开发阶段至少检查：

cpp

kernel<<<blocks, threads>>>(...);
cudaError_t launchError = cudaGetLastError();
cudaError_t runtimeError = cudaDeviceSynchronize();

生产代码应统一封装 CUDA API 错误处理，并记录设备、尺寸和调用上下文。

8. 常见初学者问题

忘记检查边界导致越界访问
频繁进行小块 Host-Device 拷贝
每个小操作都启动一个 Kernel
假设 Kernel 完成后 CPU 才继续执行
在分支中错误使用块内 Barrier
只关注线程数量，不关注访存模式
用默认 Stream 隐式同步整个流水线

8. Host 与 Device 边界

Host 代码负责：

设备发现；
内存分配；
数据准备；
Kernel 配置；
Stream/Event；
错误检查；
结果提交；
资源释放。

Device Kernel 负责大量并行线程上的局部计算。

边界越细，启动、传输和同步成本越高。

9. 线程索引

一维索引：

cpp

std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
    output[i] = transform(input[i]);
}

边界检查必须覆盖最后一个不完整 Block。

二维和三维索引要检查 Pitch、Leading Dimension 和展平方式。

整数类型需要容纳最大索引乘积。

10. Grid-Stride Loop

cpp

for (std::size_t i = blockIdx.x * blockDim.x + threadIdx.x;
     i < n;
     i += blockDim.x * gridDim.x) {
    output[i] = transform(input[i]);
}

它允许有限 Grid 覆盖大输入，并有利于线程复用。

Grid 大小仍需结合设备、Kernel 和工作量测量。

11. Block 大小

Block 大小影响：

Warp 数；
寄存器；
Shared Memory；
Occupancy；
调度；
尾部；
分支；
数据分块。

常见倍数不是固定最优值。使用 Occupancy API、编译报告和 Nsight 测量。

12. 同步范围

Block 内同步只协调同一 Block。

cpp

__syncthreads();

所有活跃线程必须以一致控制流到达 Barrier。分支中部分线程跳过会导致错误或死锁。

跨 Block 同步通常需要 Kernel 边界或专用 Cooperative 模式。

13. 原子操作

原子适合保护单个更新。

高竞争地址会串行化。

常见改善：

Block 局部归约；
Warp 聚合；
分片；
Histogram 私有化；
两阶段合并。

原子顺序与浮点结果需要单独验证。

14. Stream

Stream 定义设备操作顺序。

同一 Stream 内按序执行。不同 Stream 在资源和依赖允许时可能并发。

text

stream 0: H2D A -> Kernel A -> D2H A
stream 1: H2D B -> Kernel B -> D2H B

异步不等于必然重叠。需要 Pinned Memory、独立引擎、资源和正确依赖。

15. Event

CUDA Event 可以：

建立 Stream 依赖；
测量设备时间；
标记阶段；
查询完成。

Host 墙钟与 Device Event 回答不同问题。端到端报告两者。

16. 默认 Stream

默认 Stream 语义受编译和运行配置影响。

不要依赖隐式全局同步来保证正确性。显式表达 Stream 和 Event 依赖。

17. 异步错误

Kernel 启动错误与执行错误可能在不同时间出现。

cpp

kernel<<<grid, block, 0, stream>>>(...);
check(cudaGetLastError());
check(cudaStreamSynchronize(stream));

开发阶段在关键边界同步有助定位。生产路径减少无必要同步，但保留错误协议。

18. RAII

设备内存、Stream 和 Event 应使用 RAII 包装。

异常、早退和取消路径也必须释放资源。

销毁前确认仍在途操作不再引用对应 Buffer。

19. 多设备

多 GPU 代码需要显式设备上下文。

cpp

check(cudaSetDevice(deviceId));

资源属于创建它的设备。

P2P、拓扑、NUMA 和集合通信影响扩展。

20. CUDA Graph

CUDA Graph 可以把重复启动序列捕获为图，减少 Host 启动开销。

适合：

迭代结构稳定；
Kernel 很多；
启动开销明显；
参数可更新。

动态控制和资源生命周期会增加图管理复杂度。

21. Cooperative Groups

Cooperative Groups 提供更明确的线程组抽象。

使用时仍受设备能力、Grid 驻留和启动约束。

不要把跨 Grid 协作当成通用替代 Kernel 边界。

22. Profile 流程

text

Nsight Systems
  -> find copy, idle, launch and dominant kernels
Nsight Compute
  -> inspect selected kernel metrics
change
  -> return to system timeline and benchmark

Kernel 级提升必须回到端到端验证。

23. 测试

CUDA 测试覆盖：

n=0；
小于 Warp；
非 Block 倍数；
大输入；
非法参数；
OOM；
多 Stream；
多设备；
取消；
CPU 参考；
数值容差；
Sanitizer。

24. 完成标准

应能：

正确计算索引；
解释 Grid/Block/Warp；
选择同步范围；
管理 Stream/Event；
处理异步错误；
保护资源生命周期；
建立 CPU 参考；
用 Nsight 定位；
验证端到端收益。

核心总结

CUDA 用 Grid、Block、Thread 描述大规模并行工作。
Block 是共享内存与同步的基本协作范围。
Kernel 启动、显存分配和数据传输都是成本。
Stream 让传输与计算有机会重叠。
正确性、错误检查和边界处理先于性能优化。

下一篇：[[07-parallel-patterns]]

CUDA 编程模型：Thread、Block、Grid 与内存协作 / CUDA Threads, Blocks, Grids, and Cooperative Memory Access ​

1. Host 与 Device ​

2. 第一个 Kernel ​

3. Grid、Block 与 Thread ​

4. 线程块怎样映射到硬件 ​

5. Shared Memory 协作 ​

6. 异步执行与 Stream ​

7. 错误检查 ​

8. 常见初学者问题 ​

8. Host 与 Device 边界 ​

9. 线程索引 ​

10. Grid-Stride Loop ​

11. Block 大小 ​

12. 同步范围 ​

13. 原子操作 ​

14. Stream ​

15. Event ​

16. 默认 Stream ​

17. 异步错误 ​

18. RAII ​

19. 多设备 ​

20. CUDA Graph ​

21. Cooperative Groups ​

22. Profile 流程 ​

23. 测试 ​

24. 完成标准 ​

核心总结 ​

CUDA 编程模型：Thread、Block、Grid 与内存协作 / CUDA Threads, Blocks, Grids, and Cooperative Memory Access

1. Host 与 Device

2. 第一个 Kernel

3. Grid、Block 与 Thread

4. 线程块怎样映射到硬件

5. Shared Memory 协作

6. 异步执行与 Stream

7. 错误检查

8. 常见初学者问题

8. Host 与 Device 边界

9. 线程索引

10. Grid-Stride Loop

11. Block 大小

12. 同步范围

13. 原子操作

14. Stream

15. Event

16. 默认 Stream

17. 异步错误

18. RAII

19. 多设备

20. CUDA Graph

21. Cooperative Groups

22. Profile 流程

23. 测试

24. 完成标准

核心总结