附录 A. 速查表
A.1 公式
| 量 | 公式 | 备注 |
|---|---|---|
| 算术强度 | \(I=\text{FLOPs}/\text{Bytes}\) | FLOP/B |
| roofline | \(T=\max\left(\frac{\text{FLOPs}}{P},\frac{\text{Bytes}}{BW}\right)\) | 取大者 |
| 屋脊点 | \(I^{*}=P/BW\) | H100 FP32 约 20，TF32 约 148 |
| 在途字节 | \(N_{\text{inflight}}=BW\cdot t_{\text{lat}}\) | H100 约 2 MB |
| 占用率 | \(\dfrac{\text{block/SM}\times(\text{blockDim}/32)}{64}\) | 上限 64 warp/SM |
| 有效带宽 | \(BW_{\text{eff}}=\text{Bytes}/t\) | 和峰值比 |
| Amdahl | \(S=\dfrac{1}{(1-f)+f/p}\) | 串行段决定上限 |
| work / span | \(T_p\ge\max(W/p,\ S)\) | GPU 上 span 更常成为瓶颈 |
| 平方 tile 强度 | \(I_{\text{tile}}=T/4\) | \(T=32\to8\)，\(T=128\to32\) |
| 矩形 tile 强度 | \(I_{\text{tile}}=\dfrac{BM\cdot BN}{2(BM+BN)}\) | DRAM 侧 |
A.2 启动配置
| 项 | 取值 | 说明 |
|---|---|---|
| block 大小 | 128–256，且是 32 的倍数 | 上限 1024；尾 warp 不要有空 lane |
| grid 大小 | \(\lceil N/\text{block}\rceil\)，或 perSM × SM 数 | 后者配合 grid-stride |
| 动态 shared | 第三启动参数传字节数 | 超过 48 KB 需 cudaFuncSetAttribute |
| 寄存器约束 | __launch_bounds__(256, 4) | 比全局 -maxrregcount 精准 |
| 共享内存转置 | float tile[32][33] | +1 消 bank conflict |
| block 形状 | dim3(32, 8) | 二维索引时 x 最快变化 |
A.3 API 速查
| 类别 | 常用调用 |
|---|---|
| 设备 | cudaGetDeviceCount、cudaGetDeviceProperties、cudaDeviceGetAttribute |
| 分配 | cudaMalloc、cudaMallocManaged、cudaMallocAsync、cudaFree |
| 拷贝 | cudaMemcpy、cudaMemcpyAsync、cudaMemset |
| 锁页 | cudaHostAlloc、cudaHostRegister、cudaFreeHost |
| 流 | cudaStreamCreate、cudaStreamSynchronize、cudaStreamDestroy |
| 事件 | cudaEventRecord、cudaEventSynchronize、cudaEventElapsedTime |
| 错误 | cudaGetLastError、cudaDeviceSynchronize、cudaGetErrorString |
| 占用率 | cudaOccupancyMaxActiveBlocksPerMultiprocessor |
| 属性 | cudaFuncSetAttribute（动态 shared 上限） |
| 同步 | __syncthreads、__syncwarp、__threadfence、__threadfence_system |
| 原子 | atomicAdd/Sub/Min/Max/Exch/CAS/And/Or/Xor |
| warp | __shfl_sync、__shfl_down_sync、__ballot_sync、__any_sync、__all_sync、__popc、__ffs、__match_any_sync |
| 调试 | compute-sanitizer --tool memcheck|racecheck|initcheck、ncu、nsys |
A.4 症状与原因
| 症状 | 先查什么 |
|---|---|
| 比 CPU 还慢 | 传输/启动占比；数据是否常驻显存（0.5、6.1） |
| 只有峰值带宽的 10%–20% | 跨步访问、AoS、每线程 4 B（2.3） |
| 某个 kernel 慢 8 倍 | warp 内地址跨 128 B，sector 浪费（2.3） |
| 某个 kernel 慢 32 倍 | shared 32 路 bank conflict（2.5） |
| 结果随机、偶发错误 | 缺 __syncthreads、数据竞争（3.1、3.6） |
| 挂死 | 发散的 __syncthreads、跨 block 自旋（3.1、3.3） |
| 结果全 0 / 垃圾 | 没拷贝、没初始化、没同步就读取（6.8） |
| launch failed: invalid configuration | block > 1024、shared > 上限、grid 超限 |
| illegal memory access | 越界、未对齐的 float4、用错设备指针 |
| 第一次慢、后面快 | warmup 不足、时钟爬升（7.1） |
| 算力利用率只有 10% | 分支发散、依赖链长、没用 FMA/Tensor Core（1.2、5.5） |
| 占用率上不去 | 寄存器或 shared 太多、block 太小（1.5） |
A.5 术语
| 术语 | 含义 |
|---|---|
| SM | Streaming Multiprocessor，GPU 的计算单元，内含多个 warp 调度器 |
| warp | 32 个 lane，GPU 的调度单位 |
| lane | warp 里的一个执行通道，对应一个“线程” |
| block | 线程块，同步与 shared memory 的边界 |
| grid | 一次 kernel 启动的所有 block |
| occupancy | 常驻 warp 数与硬件上限之比 |
| coalescing | 一个 warp 的地址被合并成尽量少的 32 B sector |
| bank conflict | shared memory 中多个 lane 访问同 bank 的不同地址而串行 |
| halo | stencil 分块时需要额外加载的邻域边界 |
| roofline | 用算力和带宽两条斜线给出性能上界的模型 |
| WMMA | Warp Matrix Multiply-Accumulate，Tensor Core 的 warp 级接口 |
| TF32 | 19 位浮点（10 位尾数），Tensor Core 的 FP32 替代精度 |
| pinned memory | 被锁页、可被 DMA 直接访问的主机内存 |
| UVA | 统一虚拟地址，主机与设备指针在同一地址空间 |