GPU（CUDA）性能分析 — 到底是什么限制了你的 kernel？
Source: Dev.to
抱歉，我无法直接访问外部链接的内容。如果您能提供需要翻译的具体文本，我将很乐意为您翻译成简体中文，并保留原有的格式、代码块和链接。
介绍
在上一篇文章中，我介绍了 GPU Flight，一个轻量级的 CUDA 可观测性工具，充当 GPU 的飞行记录仪。它收集系统指标、设备能力以及每个 kernel 事件。
今天我们聚焦 GPU Flight 捕获的一个特定指标：占用率。占用率是 GPU 性能的关键指标，但常常被误解。
什么是 Occupancy？
GPU 被组织为流式多处理器（Streaming Multiprocessors，简称 SM）。每个 SM 可以同时运行许多线程——不是像 CPU 那样通过上下文切换，而是实际并行执行。SM 上的调度单元是 warp：一组 32 条线程，它们以锁步方式执行相同的指令。
一个 SM 有固定的 warp 配额（例如，在典型的 Ampere GPU 上为 48 个 warp）。
如果你以 256 线程（每块 8 个 warp）的块来启动 kernel，SM 最多可以同时容纳 6 块，以填满这 48 个 warp 槽。任何消耗额外资源的因素——寄存器、共享内存——都可能减少能够容纳的块数，从而导致部分 warp 槽空闲。
[ \text{occupancy} = \frac{\text{active warps}}{\text{maximum warps per SM}} ]
- 1.0 → 每个 warp 槽都被使用。
- 0.5 → 只利用了 SM 计算能力的一半。
GPU Flight 如何测量占用率
GPU Flight 会在每次 kernel 启动时自动记录占用率——无需修改代码。只需将 enableKernelDetails: true 打开，相关信息就会出现在日志中：
{
"type": "kernel_event",
"name": "_Z18block_reduce_naivePKfPfi",
"occupancy": 0.833333,
"num_regs": 16,
"static_shared_bytes": 16384,
"dyn_shared_bytes": 0,
"block": "(256,1,1)",
"grid": "(16384,1,1)",
"max_active_blocks": 5
}在内部，GPU Flight 会在 kernel 启动时调用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 来获取 max_active_blocks，随后除以 SM 的 warp 预算来计算占用率。这一过程发生在 CUPTI 回调中，对 kernel 执行 没有任何额外开销。
占用率为 0.833333 表示该 kernel 只在每个 SM 上占用了 5 个（而非最多可能的 6 个）并发 block——仍有一部分计算能力未被使用。
每资源占用率细分
要找出限制因素，GPU Flight 现在提供每资源的占用率细分，并自动识别出限制资源。
示例 Kernel（朴素块归约）
// block_reduce_naive.cu
__global__ void block_reduce_naive(const float* in, float* out, int n) {
__shared__ float smem[4096]; // 16 KB – 静态预留
int tid = threadIdx.x;
int gid = blockIdx.x * blockDim.x + tid;
// Load one element per thread into shared memory
smem[tid] = (gid 0; s >>= 1) {
if (tid >>(d_in, d_out, N);GPU Flight 会立即标记出问题：
{
"occupancy": 0.833333,
"reg_occupancy": 1.0,
"smem_occupancy": 0.833333,
"warp_occupancy": 1.0,
"block_occupancy":1.0,
"limiting_resource":"shared_mem"
}- 每个
*_occupancy字段回答：“如果仅考虑此约束，占用率会是多少？” limiting_resource表示实际的瓶颈。这里 共享内存 与整体占用率相匹配，而寄存器、warp 和块数并未成为限制因素。
为什么共享内存是罪魁祸首
__shared__ float smem[4096] 为每个块预留了 16 KB 的静态共享内存，无论 kernel 实际使用多少。每块 256 线程时，归约只会触及索引 0 … 255（1 KB），但整个 16 KB 在块的整个生命周期内都被锁定。过度的预留导致 SM 无法调度出最大数量的并发块。
Source: …
使用动态共享内存进行优化
切换到动态共享内存，其大小在启动时指定：
// block_reduce_optimized.cu
__global__ void block_reduce_optimized(const float* in, float* out, int n) {
extern __shared__ float smem[]; // size supplied at launch
int tid = threadIdx.x;
int gid = blockIdx.x * blockDim.x + tid;
smem[tid] = (gid 0; s >>= 1) {
if (tid >>(d_in, d_out, N);现在每个块的共享内存占用从 16 KB 降至 1 KB（降低了 16 倍），使得 SM 能够容纳全部 6 个并发块。
GPU Flight 验证了这一改进：
{
"occupancy": 1.0,
"limiting_resource": "warps"
}当 limiting_resource 变为 warps 时，表示已达到满占用——每个 SM 的 warp 槽都已被填满，共享内存不再是瓶颈。
要点
- Occupancy 是一个简洁的指标，揭示 SM（流式多处理器）的计算资源利用情况。
- 单一的 occupancy 数值可能掩盖根本原因；对每种资源的细分可以揭示真正的限制因素（寄存器、共享内存或块数量）。
- 使用 dynamic shared memory（或以其他方式减少静态分配）通常可以提升 occupancy 并获得更好的性能。
- GPU Flight 自动检测并报告这些问题，让你专注于修复 kernel，而无需手动计算硬件限制。