📚 课程目录第3章工作负载分析

3. 工作负载分析：识别计算密集型和内存密集型任务

3.1 核心原理：为什么先分类再调优

在多内核GPU环境中，性能瓶颈往往不是“GPU不够快”，而是工作负载特征与硬件资源不匹配。计算密集型（Compute-bound）任务受限于流处理器（SM）的算术吞吐量，而内存密集型（Memory-bound）任务受限于显存带宽或缓存命中率。若不加区分地统一调优，可能把精力浪费在错误的资源上——例如为内存密集型任务增加SM频率毫无意义，因为瓶颈在数据搬运路径上。

一个实用的分类维度是算术强度（Arithmetic Intensity, AI），定义为“每字节数据搬运对应的浮点运算次数”。设某内核执行了 F 次浮点运算，从显存读取/写入了 B 字节数据，则 AI = F / B。GPU的硬件拐点（Roofline模型中的ridge point）由峰值算力与显存带宽的比值决定：

ridge_point = 峰值FLOPS / 显存带宽 (单位: FLOP/Byte)
例如 A100: 19.5 TFLOPS(FP32) / 1.6 TB/s ≈ 12.2 FLOP/Byte

当 AI < ridge_point 时，内核为内存密集型；当 AI > ridge_point 时，为计算密集型。这一判断决定了后续优化方向：前者应聚焦数据复用、合并访问、压缩；后者应聚焦指令级并行、降低分支发散、使用更快的数学函数。

快速判断经验：若内核中每个线程仅读取1~2个浮点数就进行大量运算（如矩阵乘法分块后），通常是计算密集型；若每个线程需要连续读取数组元素且运算简单（如向量加法、归约），则大概率是内存密集型。

3.2 方法流程：使用NVIDIA Nsight进行瓶颈定位

NVIDIA Nsight Compute（ncu）是面向CUDA内核的细粒度性能分析器，可提供SM占用率、内存吞吐、指令混合等硬件计数器数据。推荐的分析流程分为四步：

粗筛（Nsight Systems）：先使用 nsys 抓取时间线，确认GPU空闲率、内核启动间隙、数据拷贝与内核执行是否重叠。若GPU利用率低于80%，问题可能在CPU端或数据传输。
细测（Nsight Compute）：对目标内核运行 ncu --set full，获取Roofline图、SM利用率、内存吞吐等关键指标。
对照Roofline：在Nsight Compute的Roofline视图中，观察内核所在位置。若落在“内存屋顶”下方，说明未充分利用带宽；若落在“计算屋顶”下方，说明算术流水线有停顿。
定位具体原因：结合Memory Workload Analysis和Scheduler Statistics，判断是未合并访问、缓存未命中、还是warp调度延迟。

典型命令示例如下：

# 1. 时间线粗筛
nsys profile --stats=true -o myapp ./myapp

# 2. 内核级细测（只分析名为matmul_kernel的内核）
ncu --kernel-name matmul_kernel --set full \
    --section MemoryWorkloadAnalysis \
    --section SchedulerStats \
    --section WarpStateStats \
    -o matmul_report ./myapp

分析输出中需要重点关注的指标包括：

指标	含义	瓶颈指示
Memory Throughput (%)	显存带宽利用率	>80% 且计算利用率低 → 内存密集型
SM Busy (%)	SM有活动warp的时间占比	<60% 可能存在调度停顿或依赖延迟
L1/TEX Cache Hit Rate	L1缓存命中率	<50% 且访问模式差 → 优化数据复用
Warp Stall Reasons	warp停顿原因分布	Long Scoreboard → 等待显存；Wait → 同步开销
Executed IPC (Instructions Per Cycle)	每周期执行指令数	远低于理论值 → 指令级并行不足

3.3 案例：向量求和与矩阵乘法的对比分析

我们以两个典型内核为例，展示如何通过Nsight数据区分工作负载类型。

案例A：向量求和（SAXPY）——每个线程处理一个元素，运算仅一次乘加，但需读取x和y两个数组并写回y。

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

Nsight报告显示：Memory Throughput ≈ 92%，SM Busy ≈ 35%，L1命中率约30%。Roofline图中该内核紧贴“内存屋顶”下方，算术强度AI≈0.5 FLOP/Byte。结论：典型内存密集型。优化方向：使用向量化加载（float4）、确保合并访问、考虑使用共享内存减少全局访问（若数据可复用）。

案例B：分块矩阵乘法（SGEMM）——每个线程计算输出矩阵的一个元素，需读取A的一行和B的一列，但通过分块（如16×16 tile）将数据缓存到共享内存。

__global__ void sgemm_v2(float* A, float* B, float* C, int M, int N, int K) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    // ... 分块加载与乘累加循环
}

Nsight报告显示：SM Busy ≈ 85%，Memory Throughput ≈ 45%，算术强度AI≈25 FLOP/Byte。Roofline图中该内核接近“计算屋顶”。结论：计算密集型。优化方向：提高指令级并行（ILP）、使用__ldg()只读缓存、调整block大小以提升占用率。

对比结论：两个内核在Nsight中的表现截然不同。若对SAXPY盲目增加线程数或使用更复杂的缓存策略，收益甚微；而对SGEMM若只优化内存访问而不关注计算流水线，同样无法突破性能天花板。先分类，再对症下药，是性能调优的第一原则。

3.4 误区与复盘：常见分析陷阱

在实际项目中，工程师常陷入以下误区，导致分析结论偏差：

只看GPU利用率，忽略具体瓶颈：GPU利用率高（如95%）不代表内核高效。若SM忙碌但大部分时间在等待显存返回（stall），实际计算吞吐很低。必须结合IPC和Stall Reasons判断。
忽略数据传输开销：当PCIe传输时间占比超过内核执行时间时，即使内核本身优化得再好，整体性能仍受限于I/O。此时应使用nsys检查Memcpy HtoD和DtoH是否与内核重叠。
将L2缓存命中率当作全局内存效率：L2命中率高可能只是数据局部性好，但若L1命中率低，仍会产生大量L2访问流量。应分层查看L1、L2、DRAM各级吞吐。
忽略工作负载的动态变化：同一内核在不同输入规模下可能从计算密集型变为内存密集型。例如小矩阵乘法（如32×32）因数据可完全放入共享内存，AI极高；而大矩阵（如4096×4096）因分块后仍需频繁访问全局内存，AI下降。因此每次调优都应基于实际运行配置重新测量。

复盘清单（每次分析后自查）：

是否已用Roofline模型确认内核位于哪个屋顶下方？
Memory Throughput与SM Busy的差距是否超过20%？若超过，瓶颈在何处？
Warp Stall Reason中占比最高的原因是什么？是否与内存延迟或同步有关？
是否检查过内核启动参数（block大小、grid大小）对占用率的影响？
数据传输与内核执行是否重叠？若没有，是否使用了CUDA Stream？
是否对比过不同编译器优化级别（-O2 vs -O3）下的性能差异？