CUDA 线程层次、内存层次与执行配置
CUDA 线程层次、内存层次与执行配置
| 层次 / 存储 | 作用 | 典型配置或容量 | 注意事项 |
|---|---|---|---|
| thread | 最小执行单元 | 由 block 维度决定 | 索引必须显式做边界检查，越界访问不一定报错 |
| warp | 调度与执行的最小单位 | 固定 32 个线程 | 同 warp 分支发散会被串行化，尽量按 warp 对齐分支 |
| block | 协作单元，可共享 shared memory | 常用 128 到 512 线程 | block 内可 __syncthreads()，block 间不可同步 |
| grid | 一次 kernel 启动的全部 block | 可达数百万 block | block 的调度顺序没有保证，不要依赖跨 block 顺序 |
| register | 最快的存储，线程私有 | 每线程几十到 255 个 | 寄存器用量过高会拉低 occupancy，可用 __launch_bounds__ 限制 |
| shared memory | block 内共享的片上存储 | 每 block 数十 KB，可配置 | 按 32 个 bank 划分，同 bank 不同地址会 bank conflict |
| global memory | 显存，容量最大、延迟最高 | 数 GB 到数十 GB | 务必合并访存，避免大跨步访问 |
| constant / texture | 只读缓存路径 | 常量区通常 64 KB | 适合整个 warp 读同一个地址的场景 |
执行配置的写法是 kernel<<<gridDim, blockDim, sharedBytes, stream>>>(args)，四个参数依次是 grid 维度、block 维度、动态 shared memory 字节数、所属 stream。gridDim 与 blockDim 都可以是一维到三维，最常见的是按元素总数除以每线程处理的元素数算出 grid，再在 kernel 里补边界检查。
合并访存是性能第一原则：同一个 warp 的 32 个线程访问连续且对齐的地址时，硬件会把请求合并成尽可能少的显存事务。矩阵按列访问、结构体数组（AoS）布局、跨步 stride 大于 1 的访问，都会把一次事务拆成很多次，带宽利用率断崖式下降。常用手法是向量化（用 float4 一次搬 16 字节）以及用 shared memory 做重排。
bank conflict：shared memory 按 4 字节宽度划分成 32 个 bank，同一个 warp 内有两个线程访问同一 bank 的不同地址就要排队。经典场景是矩阵转置时按列读，给每行加一个 padding（比如声明成 [32][33]）就能打散。同一地址的访问属于广播，不冲突，这一点常被误解。
stream 与异步：用 cudaMemcpyAsync 配合 stream 让拷贝与计算重叠，前提是 host 侧内存用 cudaMallocHost 申请为页锁定内存，否则异步退化成同步；多条 stream 之间用 event 建立依赖，避免一条 stream 的同步点卡住整个设备。
