2、计算瓶颈剖析：GPU利用率低的原因分析、算子瓶颈识别、CUDA Kernel优化基础

好，咱们直接切入正题。很多同学跑模型，一看GPU利用率才30%，心里就发毛。我刚开始做训练优化那会儿，也经常对着nvidia-smi发呆——明明显存占满了，算力就是上不去。这到底是谁在拖后腿？

说白了，GPU利用率低，通常就三个原因：数据加载太慢、算子效率太低、Kernel启动开销太大。咱们一个一个拆。

2.1 GPU利用率低：三大典型原因

先看数据加载。你想想看，GPU算得飞快，但CPU来不及喂数据，那GPU就只能干等。这种现象叫数据饥饿。我在项目中遇到过，用PyTorch的DataLoader默认设置，单进程加载图片，GPU利用率直接掉到20%。后来加了num_workers=8，配合prefetch_factor=2，才把利用率拉回90%以上。

再看算子效率。有些算子天生就慢，比如逐元素操作（element-wise）。你写个y = x + 1，看起来简单，但每个元素都要启动一个Kernel。我见过一个项目，把LayerNorm里的逐元素操作合并成融合Kernel，速度直接快了3倍。

最后是Kernel启动开销。每次调用CUDA Kernel，都有固定的启动延迟。如果你的模型里全是小算子，那大部分时间都花在启动上了，而不是真正在算。嗯，这里要注意：小算子多，不如大算子少。

瓶颈类型	典型表现	排查工具
数据加载	GPU利用率波动大，显存占用高但算力低	nsys profile、DataLoader耗时统计
算子效率	某个算子耗时占比异常高	PyTorch Profiler、NVIDIA Nsight
Kernel启动	小算子密集，GPU空闲时间片多	nsys timeline、Kernel launch分析

2.2 算子瓶颈识别：从Profiler到火焰图

怎么找到瓶颈算子？我个人习惯先用PyTorch Profiler跑一遍。它会告诉你每个算子的耗时、调用次数、CUDA Kernel名称。我建议重点关注两个指标：CUDA time占比和Kernel launch次数。

举个例子，你看到某个算子耗时占了40%，但它的计算量其实很小。那说明什么？说明这个算子可能没有充分利用GPU的并行能力。我曾经排查过一个模型，发现一个reshape操作耗时异常高。后来发现，reshape触发了内存拷贝，而不是简单的视图变换。改成view之后，耗时直接降了80%。

还有一个工具叫NVIDIA Nsight Systems，它能展示GPU上的Kernel执行时间线。你一眼就能看出哪些Kernel是串行的，哪些是并行的。我习惯先看GPU busy rate，如果低于80%，说明有大量空闲时间。

2.3 CUDA Kernel优化基础：从理论到实践

好了，找到瓶颈算子之后，怎么优化？这里我讲三个最基础、也最实用的方向：内存访问模式、线程块配置、算子融合。

2.3.1 内存访问：合并访问是关键

GPU的显存带宽很高，但前提是你的访问模式要合并。什么意思？就是相邻的线程要访问相邻的内存地址。如果你让线程0访问地址0，线程1访问地址100，那带宽利用率直接掉到10%。

我在优化一个矩阵转置Kernel时，就踩过这个坑。一开始用朴素实现，带宽利用率只有15%。后来改成共享内存分块，先让线程块内合并加载到共享内存，再转置写回，带宽利用率提升到了70%。

// 不合并访问（慢）
__global__ void bad_transpose(float* in, float* out, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    out[col * n + row] = in[row * n + col];  // 写操作不连续
}

// 合并访问（快）
__global__ void good_transpose(float* in, float* out, int n) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    tile[threadIdx.y][threadIdx.x] = in[row * n + col];
    __syncthreads();
    out[col * n + row] = tile[threadIdx.x][threadIdx.y];  // 写操作连续
}

2.3.2 线程块配置：找到最佳网格大小

线程块的大小和数量，直接影响GPU的占用率。我建议遵循一个原则：让每个SM至少跑4个Warp，这样才能隐藏内存访问延迟。

怎么算？假设你的GPU有80个SM，每个SM最多2048个线程。那你至少需要80 * 4 * 32 = 10240个线程。线程块大小一般选128或256，这样每个SM能同时跑8-16个Warp。

我曾经优化过一个逐元素Kernel，一开始线程块设成32，占用率只有25%。改成256之后，占用率到了75%，速度翻了3倍。你想想看，只是改了个数字，效果就这么明显。

线程块大小	每个SM的Warp数	占用率	性能表现
32	2	25%	差
128	8	50%	中等
256	16	75%	好

2.3.3 算子融合：减少Kernel启动

最后一个技巧，也是我最常用的——算子融合。把多个小算子合并成一个Kernel，减少启动开销和中间显存读写。

举个例子，y = (x + 1) * 2，如果分开写，要启动两个Kernel，中间结果还要写回显存。融合之后，一个Kernel搞定，中间结果留在寄存器里。

// 分开写（慢）
z = x + 1;
y = z * 2;

// 融合写（快）
__global__ void fused_kernel(float* x, float* y, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        y[idx] = (x[idx] + 1.0f) * 2.0f;  // 一次计算，无中间显存
    }
}

嗯，这里要注意：不是所有算子都能融合。比如涉及跨线程通信的（像Softmax），融合起来就比较复杂。但像逐元素操作、激活函数、简单的数学运算，几乎都能融合。

好了，这一章的内容就到这里。记住，优化GPU利用率不是一蹴而就的事，需要你反复profile、分析、改代码。但只要你掌握了上面这些方法，大部分瓶颈都能找到并解决。