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 隐式同步整个流水线

历史 CUDA 阅读补充

本页保留课程合并前的 CUDA 编程模型说明。

当前版本补充了 Stream、Event、异步错误、RAII、多设备和 CUDA Graph。

Host 职责

Host 负责设备、内存、启动、同步、错误和资源。

Host/Device 边界是性能和正确性边界。

Device 职责

Kernel 在线程层次上执行局部工作。

Kernel 不应访问越界数据或依赖未表达的跨 Block 同步。

索引

cpp

auto i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) { ... }

二维和三维需要明确展平、Stride 和边界。

Grid-Stride

Grid-Stride Loop 让有限线程覆盖大数据。

步长是 blockDim * gridDim。

Block

Block 大小影响 Warp、寄存器、Shared Memory 和 Occupancy。

常用值只是起点，最终靠 Profile。

Warp

Warp 以相同指令执行线程。

分支发散让不同路径分阶段执行。

同步

Block Barrier 只覆盖同一 Block。

所有参与线程必须一致到达。

跨 Block 通常使用 Kernel 边界。

原子

原子保证单个更新，但高竞争会串行。

局部聚合再全局合并可以减少原子。

Stream

同一 Stream 有序。不同 Stream 可能并发。

异步 API 不保证立即完成或必然重叠。

Event

Event 用于依赖、计时和完成查询。

设备时间和 Host 墙钟都应记录。

传输

H2D/D2H 必须计入端到端。

Pinned Memory 可以支持异步传输，但需要预算。

默认 Stream

默认 Stream 语义可能与配置相关。

使用显式 Stream/Event 表达关键依赖。

错误检查

Kernel 启动后检查 Launch Error。

在同步点检查执行错误。

错误必须进入任务终态和日志。

资源生命周期

Buffer、Stream、Event 使用 RAII。

释放前确认在途操作已经完成。

多 GPU

资源与创建设备关联。

处理 Device ID、P2P、拓扑、NUMA 和负载均衡。

CUDA Graph

重复而稳定的启动序列可用 Graph 降低 Host 开销。

动态图和资源更新增加复杂度。

调试

使用：

CPU 参考；
小输入；
同步定位；
Compute Sanitizer；
Nsight Systems；
Nsight Compute；
错误宏；
数值容差。

测试边界

空输入；
小输入；
非 Block 倍数；
大索引；
OOM；
多 Stream；
多设备；
取消；
错误注入。

历史版本边界

CUDA Runtime、设备能力和工具指标持续变化。

旧 API、Block 建议和性能阈值需要在目标 Toolkit 与 GPU 核对。

复现记录

保存：

CUDA Toolkit；
Driver；
GPU；
Compute Capability；
编译选项；
输入；
启动配置；
Profile 报告；
正确性。

当前课程映射

当前文章新增：

Host/Device 职责；
Grid-Stride；
Stream/Event；
异步错误；
RAII；
多设备；
CUDA Graph；
测试矩阵。

归档验收清单

Kernel 围栏完整；
线程索引说明存在；
同步范围明确；
传输计入性能；
异步错误有说明；
工具版本有边界；
CPU 参考可用；
输入规模记录；
数值容差记录；
设备与驱动记录；
历史命令可追溯；
当前文章入口明确。

历史示例如果无法在新 Toolkit 编译，应保留原意并提供迁移说明。不要只为通过编译删除错误检查或同步。重新优化时从当前端到端基线开始，而不是复用旧 GPU 的结论。归档代码的输出也必须经过相同正确性校验。性能提升不能以跳过错误检查为代价。所有结论都绑定记录的设备、驱动和 Toolkit。

核心总结

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

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

原体系结构与硬件 / Original Architecture

原并行计算 / Original Parallel Computing

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. 常见初学者问题

历史 CUDA 阅读补充

Host 职责

Device 职责

索引

Grid-Stride

Block

Warp

同步

原子

Stream

Event

传输

默认 Stream

错误检查

资源生命周期

多 GPU

CUDA Graph

调试

测试边界

历史版本边界

复现记录

当前课程映射

归档验收清单

核心总结

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. 常见初学者问题 ​

历史 CUDA 阅读补充 ​

Host 职责 ​

Device 职责 ​

索引 ​

Grid-Stride ​

Block ​

Warp ​

同步 ​

原子 ​

Stream ​

Event ​

传输 ​

默认 Stream ​

错误检查 ​

资源生命周期 ​

多 GPU ​

CUDA Graph ​

调试 ​

测试边界 ​

历史版本边界 ​

复现记录 ​

当前课程映射 ​

归档验收清单 ​

核心总结 ​

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. 常见初学者问题

历史 CUDA 阅读补充

Host 职责

Device 职责

索引

Grid-Stride

Block

Warp

同步

原子

Stream

Event

传输

默认 Stream

错误检查

资源生命周期

多 GPU

CUDA Graph

调试

测试边界

历史版本边界

复现记录

当前课程映射

归档验收清单

核心总结