第十九课：编译器如何优化 CUDA Kernel

导语：这一课解释 CUDA 编译器如何从源代码生成更高效的 GPU 指令，以及这些优化怎样影响寄存器、分支、访存和 Occupancy。

CUDA 程序最终执行的不是源代码，而是编译器生成的 SASS 指令。

因此，源代码中写了多少行、多少个变量，并不等于 GPU 实际执行多少条指令。

编译过程可以简化为：

图示：第十九课：编译器如何优化 CUDA Kernel

编译器优化的目标通常是：

• 减少指令数量；
• 减少内存访问；
• 提高指令级并行性；
• 使用更高效的硬件指令；
• 降低分支和函数调用开销；
• 合理分配寄存器；
• 根据目标 GPU 架构选择指令。

但优化并不一定始终提高性能。某些优化可能增加寄存器压力、降低 Occupancy，甚至改变浮点计算结果。

---

一、编译器优化的基本原则

编译器通常遵循“可观察行为不变”的原则。

例如：

int a = 3;
int b = a * 4;
output[i] = b;

编译器没有必要真的执行一次乘法。

它可以直接生成：

output[i] = 12;

因为两段程序的可观察结果相同。

但对于浮点数、并行内存访问、原子操作和 volatile 数据，“行为是否等价”会更加复杂。

---

二、常量折叠 Constant Folding

如果表达式的操作数在编译时已知，编译器可以直接计算结果。

源代码：

int x = 16 * 32;

编译后可能直接变为：

int x = 512;

例如卷积核尺寸固定：

constexpr int K = 3;
int kernel_size = K * K;

编译器可以直接得到：

kernel_size = 9

不需要在 GPU 运行时执行乘法。

---

三、常量传播 Constant Propagation

假设：

const int stride = 2;

int index = threadIdx.x * stride;

编译器知道 stride 永远是2，可以把它传播到所有使用位置：

int index = threadIdx.x * 2;

之后还可能继续优化为移位：

int index = threadIdx.x << 1;

常量传播可以使后续优化更容易进行。

---

四、死代码消除 Dead Code Elimination

如果某段计算结果永远不会被使用，编译器可以把它完全删除。

例如：

float x = A[i];
float y = x * 2.0f;

C[i] = 1.0f;

因为 y 没有参与最终结果，编译器可能删除：

float x = A[i];
float y = x * 2.0f;

甚至对 A[i] 的 Global Memory Load 也可能一起消失。

最终只剩：

C[i] = 1.0f;

---

为什么微基准测试容易被错误优化

假设测试乘法速度：

float x = input[i];

for (int k = 0; k < 1000; ++k) {
    x = x * 1.01f;
}

如果最后没有使用 x：

// 没有输出

编译器可能把整个循环删除。

因此测试时必须让结果产生可观察影响，例如：

output[i] = x;

但即使写入输出，也要检查编译器是否进行了其他代数简化。

---

五、死存储消除 Dead Store Elimination

假设：

C[i] = 1.0f;
C[i] = 2.0f;

如果两次写入之间没有其他线程或设备能够观察第一次写入，编译器可能删除：

C[i] = 1.0f;

只保留最后一次：

C[i] = 2.0f;

但如果使用：

• volatile；
• 原子操作；
• 线程同步；
• 外部可见副作用；

编译器不能随意消除。

---

六、公共子表达式消除

假设：

float a = x * y + z;
float b = x * y + w;

其中：

x * y

被计算两次。

编译器可能改成：

float temp = x * y;
float a = temp + z;
float b = temp + w;

这样可以减少一次乘法。

但代价是 temp 需要在寄存器中保持更长时间，可能增加寄存器压力。

所以减少指令和减少寄存器使用并不总是同时实现。

---

七、函数内联 Function Inlining

假设有一个设备函数：

__device__ float relu(float x)
{
    return x > 0.0f ? x : 0.0f;
}

Kernel 中调用：

y = relu(x);

如果不内联，理论上需要：

准备参数
跳转到函数
执行函数
返回调用位置

编译器通常会把函数体直接放到调用位置：

y = x > 0.0f ? x : 0.0f;

这叫函数内联。

---

内联的优点

• 消除函数调用和返回开销；
• 让编译器看到完整上下文；
• 促进常量传播；
• 促进死代码消除；
• 促进指令融合；
• 有利于跨函数优化。

例如：

__device__ float multiply(float x, float scale)
{
    return x * scale;
}

调用：

float y = multiply(x, 2.0f);

内联后，编译器知道 scale=2.0f，可以进一步优化。

---

内联的缺点

如果一个很大的函数在很多位置被内联，会导致：

代码体积增大
    ↓
指令缓存压力增加
    ↓
寄存器活跃范围变长
    ↓
寄存器使用量增加

因此内联不是越多越好。

CUDA 中可以使用：

__forceinline__

倾向于强制内联，也可以使用：

__noinline__

阻止内联。

但这类指示应在测量后使用，不应机械添加。

---

八、循环展开 Loop Unrolling

假设：

float sum = 0.0f;

for (int k = 0; k < 4; ++k) {
    sum += x[k] * w[k];
}

循环每次需要：

• 增加循环变量；
• 判断循环条件；
• 执行分支；
• 计算数组索引。

编译器可能展开为：

sum += x[0] * w[0];
sum += x[1] * w[1];
sum += x[2] * w[2];
sum += x[3] * w[3];

这样就消除了循环控制指令。

---

循环展开的优势

第一，减少循环分支和计数指令。

第二，让编译器看见每次迭代的具体索引。

第三，便于常量传播和地址计算简化。

第四，可以增加独立指令数量，提高 ILP。

第五，可能便于生成向量化 Load/Store。

---

#pragma unroll

可以提示编译器展开循环：

#pragma unroll
for (int k = 0; k < 8; ++k) {
    sum += a[k] * b[k];
}

也可以指定展开因子：

#pragma unroll 4
for (int k = 0; k < 16; ++k) {
    ...
}

或者要求不展开：

#pragma unroll 1
for (...) {
    ...
}

---

九、循环展开为什么可能降低性能

假设循环展开后，同时出现大量中间变量：

a0、a1、a2、a3……
b0、b1、b2、b3……
sum0、sum1、sum2、sum3……

这会增加寄存器数量。

结果可能是：

图示：九、循环展开为什么可能降低性能

更严重时，还会发生 Register Spill。

此外，过度展开会增加代码体积，使：

• 指令缓存命中率下降；
• 分支目标更远；
• 编译时间增加；
• SASS 难以调度。

因此适度展开通常有利，但完全展开大型循环可能适得其反。

---

十、指令融合与 FMA

源代码：

float y = a * b + c;

编译器通常会生成一条 FMA：

$$y=a\times b+c$$

SASS 中可能是：

FFMA

而不是分成：

FMUL
FADD

---

FMA 的优点

1. 指令数量减少

两条算术操作变成一条指令。

2. 吞吐率更高

GPU FP32 执行管线通常针对 FMA 进行优化。

3. 精度可能更好

分开的乘法和加法通常会经历两次舍入：

乘法结果舍入
    ↓
加法结果再次舍入

FMA 通常只在最终结果处舍入一次。

---

FMA 会改变数值结果

由于舍入方式不同：

a * b + c

生成 FMA 后，结果可能与严格执行乘法再执行加法略有不同。

这种差异通常很小，但如果要求逐位一致，就需要谨慎。

编译器选项可以控制是否允许 FMA 融合，但关闭通常会降低性能。

---

十一、代数变换

编译器可能将：

$$x\times2$$

改为：

$$x+x$$

或者将整数除以2改为右移。

对于整数且语义允许时：

x / 8

可能变为：

x >> 3

但对于有符号负数、溢出和舍入规则，编译器必须保持语言语义。

---

浮点代数变换更复杂

数学上：

$$(a+b)+c=a+(b+c)$$

但在浮点数中，这两个结果可能不同。

例如：

a 很大
b 与 a 符号相反
c 很小

不同加法顺序会产生不同舍入误差。

因此在严格浮点语义下，编译器不能随意重新排列所有运算。

开启快速数学选项后，编译器可能允许更多重排。

---

十二、--use_fast_math

使用：

nvcc --use_fast_math kernel.cu

会启用一组更激进的浮点优化。

可能包括：

• 使用近似特殊函数；
• 使用较快的除法和平方根；
• 允许某些低精度路径；
• 对次正规数采用 Flush-to-Zero；
• 使用更宽松的浮点语义。

例如：

sinf(x)

可能使用较快的近似实现。

---

Fast Math 的代价

• 数值精度下降；
• 极端输入误差增大；
• 与 CPU 结果不一致；
• 不同 GPU 架构结果可能略有变化；
• 次正规数可能被处理成0；
• 科学计算或训练收敛可能受到影响。

因此使用前应评估：

性能提升
与
数值误差

不能只看执行时间。

---

十三、地址计算优化

考虑：

int index = row * width + col;
float x = A[index];

每次访问都可能涉及整数乘法和加法。

如果 width 是编译时常量，例如1024：

index = row * 1024 + col;

编译器可能用移位实现：

index = (row << 10) + col;

如果多个访问具有相同基址：

A[index]
A[index + 1]
A[index + 2]

编译器还可能复用基础地址，避免每次重新计算完整地址。

---

十四、循环中的地址递增

源代码可能写成：

for (int k = 0; k < K; ++k) {
    sum += A[base + k];
}

编译器不一定每轮都重新执行：

base + k

它可能维护一个指针：

const float* ptr = A + base;

for (...) {
    sum += *ptr;
    ++ptr;
}

硬件层面可能只需要每次递增地址。

这种优化称为 Induction Variable Optimization。

---

十五、分支转谓词 Predication

源代码：

if (x > 0.0f) {
    y = x;
} else {
    y = 0.0f;
}

编译器可能不生成真正跳转，而是生成：

比较
生成 Predicate
根据 Predicate 选择结果

例如概念上：

P = x > 0
P ? y = x
!P ? y = 0

这样可以避免短分支的控制流开销。

---

什么时候编译器倾向使用谓词

通常在分支体较短时，例如：

if (condition) {
    x += 1;
}

如果分支体很大：

if (condition) {
    执行几十条复杂指令
}

即使使用谓词，被屏蔽线程也可能需要经历这些指令的发射，因此真正分支可能更合适。

编译器会根据：

• 分支体长度；
• 目标架构；
• 指令成本；
• 可预测的发散情况；

选择谓词或跳转。

---

十六、Warp 一致分支优化

假设条件只由 blockIdx.x 决定：

if (blockIdx.x < 4) {
    ...
}

同一个 Block 中所有线程看到相同条件。

因此同一个 Warp 中线程不会产生分支发散。

编译器和硬件仍会执行分支，但不会产生两条路径的 Active Mask 分裂。

可以称为：

Warp-Uniform Branch

相反：

if (threadIdx.x & 1) {
    ...
}

同一 Warp 中奇偶线程走不同路径，容易产生发散。

---

十七、Load 的消除与复用

假设：

float a = A[i];
float b = A[i];
float c = a + b;

如果编译器确定 A[i] 在两次读取之间不会被修改，它可以只加载一次：

float a = A[i];
float c = a + a;

这样减少 Global Memory Load。

但如果编译器认为：

A

可能通过其他指针被修改，就不一定能够进行这种优化。

这涉及指针别名分析。

---

十八、什么是指针别名

假设函数为：

__global__ void kernel(
    float* A,
    float* B)
{
    A[i] = 1.0f;
    float x = B[i];
}

编译器不能天然确定：

A 和 B 是否指向同一片内存

如果：

A == B

那么对 A[i] 的写入会影响随后对 B[i] 的读取。

因此编译器必须采取保守策略。

---

十九、__restrict__ 的作用

可以声明：

__global__ void kernel(
    float* __restrict__ A,
    const float* __restrict__ B)

其含义是向编译器承诺：

在该指针作用范围内，相关内存不会通过其他别名指针访问。

这样编译器可以更大胆地：

• 缓存加载值；
• 消除重复 Load；
• 重新排序部分访问；
• 将数据长期保存在寄存器；
• 进行更强的向量化。

---

错误使用 __restrict__ 的后果

如果实际上 A 和 B 指向重叠区域，却声明为 __restrict__，程序行为可能错误。

例如：

kernel<<<...>>>(ptr, ptr);

但函数声明承诺 A 和 B 不重叠。

编译器可能基于错误承诺进行优化，得到不符合预期的结果。

因此 __restrict__ 是程序员向编译器提供的保证，不是运行时检查。

---

二十、const 的作用

例如：

const float* input

表示当前代码不能通过 input 修改数据。

这可以帮助编译器理解访问语义。

但 const 不一定意味着：

• 数据永远不会被其他线程修改；
• 数据一定进入某个特定缓存；
• 编译器完全确定没有别名。

所以 const 和 __restrict__ 的作用不同。

常见组合是：

const float* __restrict__ input

表示：

• 当前函数只读取；
• 不会通过别的别名指针访问同一数据区域。

---

二十一、向量化 Load/Store

假设一个线程连续读取4个 FP32：

float a = A[i];
float b = A[i + 1];
float c = A[i + 2];
float d = A[i + 3];

如果地址和对齐条件合适，可以使用：

float4 value =
    reinterpret_cast<const float4*>(A)[i / 4];

逻辑上一次读取：

$$4\times32=128\text{ bit}$$

编译器可能生成较宽的 Load 指令。

---

向量化的可能优势

• 减少 Load/Store 指令数量；
• 减少地址计算；
• 提高单线程内存指令效率；
• 适合数据打包和搬运。

但向量化不一定减少实际显存事务数量。

Global Memory 事务仍然由整个 Warp 的地址模式决定。

所以应区分：

线程级指令数量

与：

Warp 级内存事务数量

---

向量化的对齐要求

float4 通常要求16字节对齐。

如果地址没有正确对齐：

• 可能生成更多指令；
• 可能降低性能；
• 某些低级写法甚至可能产生未定义行为。

因此不能随意把任意 float* 强制转换成 float4*。

---

二十二、标量替换 Scalar Replacement

假设一个线程使用小数组：

float values[4];

values[0] = ...;
values[1] = ...;
values[2] = ...;
values[3] = ...;

如果索引在编译时可确定，编译器可以把数组元素拆成独立标量：

value0
value1
value2
value3

并放入寄存器。

这称为 Scalar Replacement of Aggregates。

---

动态索引可能导致 Local Memory

如果访问为：

float x = values[index];

且 index 在运行时变化，编译器可能无法把它映射到固定寄存器。

于是数组可能被放入 Local Memory。

因此下面两种代码可能性能差异很大：

values[2];

和：

values[index];

即使数组只有几个元素。

检查编译报告和 SASS 才能确认是否发生 Local Memory 访问。

---

二十三、寄存器分配

PTX 使用虚拟寄存器，但物理 GPU 每个 SM 的寄存器数量有限。

ptxas 需要把大量虚拟变量映射到有限物理寄存器：

虚拟寄存器
    ↓
活跃区间分析
    ↓
复用不重叠的物理寄存器
    ↓
最终每线程寄存器数量

例如两个变量生命周期不重叠：

float a = ...;
使用 a 完成后

float b = ...;

编译器可能让 a 和 b 使用同一个物理寄存器。

---

二十四、变量生命周期为什么重要

考虑：

float a = load_a();
float b = load_b();
float c = load_c();
float d = load_d();

result = a + b + c + d;

在执行加法之前，a、b、c、d 都必须保持有效，因此同时占用寄存器。

如果改写成：

float sum = load_a();
sum += load_b();
sum += load_c();
sum += load_d();

同时活跃变量可能更少。

但第二种写法形成更长的依赖链，可能降低 ILP。

因此存在权衡：

更多并行临时变量
→ ILP 更高
→ 寄存器更多

更少临时变量
→ 寄存器更少
→ 依赖链更长

编译器需要在这两者之间进行选择。

---

二十五、Register Spill

当物理寄存器无法容纳所有活跃变量时，部分值会被存入 Local Memory。

例如 SASS 中出现类似：

STL
LDL

或者其他 Local Memory 访问指令。

流程为：

寄存器不足
    ↓
中间变量写入 Local Memory
    ↓
以后再从 Local Memory 读回

虽然 Local Memory 对每个线程逻辑私有，但物理上通常位于显存地址空间，并经过缓存。

Spill 可能造成：

• 更多内存指令；
• 更高延迟；
• L1/L2 流量增加；
• 显存带宽增加；
• 性能显著下降。

---

二十六、限制寄存器数量

可以使用：

nvcc -maxrregcount=64 kernel.cu

限制每线程最多使用的寄存器数量。

也可以使用：

__launch_bounds__(threads_per_block, min_blocks_per_sm)

向编译器提供预期启动配置。

例如：

__global__
__launch_bounds__(256, 2)
void kernel(...)
{
    ...
}

表示期望：

• 每个 Block 最多256线程；
• 每个 SM 至少能够驻留2个这样的 Block。

编译器可能据此控制寄存器使用。

---

限制寄存器为什么可能反而更慢

限制寄存器可能提高 Occupancy：

寄存器减少
    ↓
驻留 Warp 增加

但也可能产生 Spill：

寄存器减少过多
    ↓
Local Memory Load/Store 增加

最终性能可能下降。

因此寄存器限制必须实际测试，不能只看 Occupancy。

---

二十七、指令调度

假设有：

LDG R1, [A]
FADD R2, R1, R3

第二条指令依赖第一条 Load 返回的数据。

如果立即排列在一起，可能形成长等待。

编译器可能把独立指令放在中间：

LDG R1, [A]
IMAD R8, R9, R10, R11
LDG R4, [B]
FADD R6, R7, R5
FADD R2, R1, R3

目的是在等待 Load 时执行其他工作。

这叫静态指令调度。

GPU 主要依赖：

• 编译器静态调度；
• Scoreboard；
• Warp 切换；

而不是像高性能 CPU 那样使用大型乱序执行窗口。

---

二十八、软件流水化

假设循环每轮需要：

1. 加载数据；
2. 计算；
3. 写回。

朴素方式：

Load 0
等待
Compute 0
Load 1
等待
Compute 1

优化后可能变成：

Load 0
Load 1
Compute 0
Load 2
Compute 1
Load 3
Compute 2

即在计算当前数据时，提前加载后续数据。

这称为 Software Pipelining。

它可以提高 Load 与计算的重叠，但也会增加：

• 同时在途数据；
• 活跃寄存器数量；
• 代码复杂度。

---

二十九、编译器能否自动使用 Shared Memory

通常不能期待编译器自动把普通 Global Memory 数据完整地 Tiling 到 Shared Memory。

例如矩阵乘法中：

A[row][k]
B[k][col]

编译器一般不会自动推断一个复杂的 Block 级协作加载方案。

原因包括：

• 需要线程协作；
• 需要同步；
• Tile 大小与硬件资源相关；
• 数据复用模式复杂；
• 边界处理复杂。

因此 Shared Memory Tiling 往往需要：

• 程序员显式编写；
• 高级编译器框架生成；
• 使用 cuBLAS、cuDNN、CUTLASS 等库。

---

三十、编译器能否自动合并访存

编译器可以优化单线程的地址计算和 Load 指令，但 Warp 访存是否合并，主要取决于程序的数据布局和线程索引。

例如：

A[threadIdx.x]

地址天然连续。

而：

A[threadIdx.x * 1024]

地址跨度很大。

编译器一般不能在保持语义的情况下把后者自动改成连续访问。

因此合并访存主要是算法映射和数据布局问题，而不是简单编译选项问题。

---

三十一、编译期特化

假设矩阵 Tile 大小是运行时参数：

void kernel(..., int tile_size)

编译器必须生成通用代码：

运行时循环
运行时边界
动态地址计算

如果使用模板：

template<int TILE>
__global__ void kernel(...)

并实例化：

kernel<16><<<...>>>();
kernel<32><<<...>>>();

编译器知道具体 Tile 大小，可以进行：

• 完整或部分循环展开；
• 常量传播；
• 固定数组分配；
• 地址计算简化；
• 分支消除。

这叫编译期特化。

---

特化的代价

每种参数组合都会生成一份 Kernel：

TILE=16 → 一份 SASS
TILE=32 → 一份 SASS
TILE=64 → 一份 SASS

参数组合过多会造成：

• 编译时间长；
• 二进制膨胀；
• 指令缓存和部署压力；
• 测试组合增多。

高性能库常在运行时从有限数量的特化 Kernel 中选择一个。

---

三十二、分支消除

假设模板参数决定功能：

template<bool USE_BIAS>
__global__ void kernel(...)
{
    float y = ...;

    if (USE_BIAS) {
        y += bias[i];
    }

    output[i] = y;
}

实例化：

kernel<false>

时，编译器知道条件永远为假，可以完全删除 Bias 路径。

相比运行时参数：

if (use_bias)

编译期特化可以消除分支和相关 Load。

---

三十三、跨模块优化与设备链接

如果设备函数位于不同源文件，传统独立编译可能使编译器看不到完整函数体。

这会限制：

• 内联；
• 常量传播；
• 跨函数死代码消除。

Device Link-Time Optimization 可以在链接阶段对多个设备模块做进一步优化。

概念上：

文件 A 的 Device Code
文件 B 的 Device Code
          ↓
设备链接阶段
          ↓
跨模块内联与优化

这适合大型 CUDA 工程，但会增加链接时间和工具链复杂度。

---

三十四、volatile 的作用

如果声明：

volatile int* flag;

编译器通常需要保留对 flag 的实际内存访问，不能随意把它长期缓存到寄存器或删除重复访问。

例如：

while (*flag == 0) {
}

没有 volatile 时，编译器可能认为 flag 在当前线程中没有被修改，从而只加载一次。

使用 volatile 后，循环中通常会重新读取。

---

volatile 不是什么

volatile 不保证：

• 原子性；
• 线程间同步；
• 内存顺序；
• Cache 一致性；
• 多线程程序正确性。

例如：

volatile int counter;
counter++;

仍然不是原子操作。

多个线程执行时仍会发生数据竞争。

需要原子更新时，应使用：

atomicAdd

需要可见性和顺序时，还可能需要：

• __syncthreads()；
• __threadfence()；
• C++ 原子内存序；
• Cooperative Groups 等机制。

---

三十五、编译器重排与硬件重排

需要区分两种重排。

编译器重排

编译器为了提高性能，重新安排指令顺序。

硬件和内存系统行为

写缓冲、Cache、互连和内存控制器可能使其他线程观察到不同的可见顺序。

即使通过 volatile 阻止部分编译器优化，也不一定建立完整的硬件内存顺序。

因此并行同步不能仅依赖 volatile。

---

三十六、为什么优化有时会增加指令数量

某种优化可能为了减少延迟而复制计算。

例如与其保存一个值很长时间占用寄存器，编译器可能在两个位置重复计算一个简单表达式：

多执行一次整数加法

以换取：

缩短寄存器生命周期
降低寄存器压力

所以“指令更少”并不总是“性能更高”。

编译器真正考虑的是综合代价：

• 指令数；
• 延迟；
• 寄存器；
• Occupancy；
• 代码大小；
• 目标执行单元吞吐率。

---

三十七、为什么优化可能降低 Occupancy

例如循环展开前：

每线程使用32个寄存器

展开后：

每线程使用72个寄存器

虽然指令级并行性提高，但一个 SM 可驻留 Warp 数可能下降。

最终可能出现两种情况。

情况一：优化有效

单 Warp 执行效率提高很多
Occupancy 虽下降，但仍足够隐藏延迟

情况二：优化有害

Occupancy 降得太低
内存延迟无法隐藏
整体性能下降

必须通过实际性能分析判断。

---

三十八、为什么优化可能导致指令缓存问题

如果 Kernel 经过大量展开和特化，SASS 代码可能非常大。

执行时指令需要从：

Instruction Cache

读取。

代码太大可能造成：

• 指令缓存 Miss；
• 前端取指停顿；
• 分支目标频繁切换；
• Warp Scheduler 没有指令可发射。

这在大型融合 Kernel、复杂模板和过度展开中更常见。

---

三十九、Kernel Fusion 与编译器优化

原来有两个 Kernel：

Kernel A：
读取 X
计算 Y
写回 Y

Kernel B：
读取 Y
计算 Z
写回 Z

融合后：

读取 X
计算 Y
Y 保存在寄存器
继续计算 Z
写回 Z

可以消除中间显存读写。

编译器还可以跨原本两个阶段进行：

• 常量传播；
• 死代码消除；
• 指令融合；
• 中间变量寄存器化。

但融合也可能使：

• 每线程寄存器增加；
• 代码体积增加；
• Occupancy 下降；
• Block 映射不再适合两个阶段；
• 同步需求增加。

所以融合也是典型的权衡。

---

四十、调试编译为什么通常很慢

使用：

nvcc -G

生成 Device Debug 代码时，许多优化会被关闭或显著削弱。

这样有利于：

• 单步调试；
• 观察变量；
• 源代码和指令对应；
• 保留函数和变量结构。

但可能导致：

• 指令数量增加；
• 寄存器增加；
• Spill 增加；
• 性能远低于 Release 版本。

因此不能用 -G 编译版本进行正式性能测试。

通常性能分析使用优化编译，并可加入：

-lineinfo

保留行号信息，而不关闭主要优化。

---

四十一、如何查看编译器到底做了什么

第一步，查看资源使用：

nvcc -O3 \
  -arch=sm_80 \
  -Xptxas=-v \
  kernel.cu

重点观察：

Used registers
spill stores
spill loads
shared memory
stack frame

第二步，生成 PTX：

nvcc -O3 \
  -arch=compute_80 \
  -ptx kernel.cu \
  -o kernel.ptx

第三步，查看 SASS：

nvcc -O3 \
  -arch=sm_80 \
  -cubin kernel.cu \
  -o kernel.cubin

nvdisasm kernel.cubin

第四步，配合 Nsight Compute 查看运行时瓶颈。

---

四十二、一个具体优化示例

考虑简单点积：

__global__ void dot4(
    const float* A,
    const float* B,
    float* C,
    int n)
{
    int i = blockIdx.x * blockDim.x
          + threadIdx.x;

    if (i < n) {
        int base = i * 4;

        float sum = 0.0f;

        for (int k = 0; k < 4; ++k) {
            sum += A[base + k]
                 * B[base + k];
        }

        C[i] = sum;
    }
}

编译器可能进行以下优化。

1. 展开循环

sum += A[base + 0] * B[base + 0];
sum += A[base + 1] * B[base + 1];
sum += A[base + 2] * B[base + 2];
sum += A[base + 3] * B[base + 3];

2. 简化 i * 4

可能使用移位或高效整数乘加。

3. 生成 FMA

每个：

sum += a * b;

可能生成：

FFMA

4. 合并地址计算

复用 A 和 B 的基础地址。

5. 使用寄存器累加

sum 全程保存在寄存器中。

6. 可能使用向量 Load

如果程序写法、对齐和别名信息允许，可能生成较宽 Load；但不能仅凭源代码假定一定发生。

最终 SASS 与源代码的循环结构可能已经明显不同。

---

四十三、编译器不能替代算法设计

编译器可以把：

十条低效指令

优化成：

五条高效指令

但它通常无法自动把一个低计算强度算法变成高计算强度算法。

例如编译器一般不能自动完成：

图示：四十三、编译器不能替代算法设计

这需要算法、库和架构级设计。

同样，编译器也很难自动修复：

• 错误线程映射；
• 严重非合并访存；
• 不合理数据布局；
• 大量跨 Block 同步；
• 算法本身的冗余数据搬运。

所以编译器优化是在正确算法映射之上的第二层优化。

---

四十四、编译器优化的正确分析方法

不应只看某一个指标。

例如寄存器从48增加到64：

可能不好：
Occupancy 下降

也可能很好：
Spill 消失
ILP 增加
访存减少

正确分析流程是：

图示：四十四、编译器优化的正确分析方法

GPU 优化必须同时关注：

• 正确性；
• 数值精度；
• 指令；
• 资源占用；
• 实际执行时间。

---

本课核心结论

第一，编译器会通过常量折叠、常量传播和死代码消除减少无效计算。

第二，函数内联可以消除调用开销并促进跨函数优化，但可能增加代码体积和寄存器压力。

第三，循环展开可以减少循环控制、提高 ILP，但可能降低 Occupancy 或造成指令缓存压力。

第四，乘法和加法通常会融合成 FMA，减少指令并改善数值精度，但结果可能与分步计算略有差异。

第五，快速数学可以提高特殊函数、除法等操作速度，但会放宽浮点精度要求。

第六，编译器会优化地址计算、循环变量和公共子表达式。

第七，短分支可能被转换为谓词执行，长分支通常更适合真正跳转。

第八，__restrict__ 可以帮助编译器进行别名分析，但错误承诺会导致错误结果。

第九，向量 Load/Store 可以减少单线程指令数，但仍需保证对齐，并不自动保证 Warp 级访存合并。

第十，动态索引的小数组可能进入 Local Memory，而固定索引数组更容易被拆到寄存器中。

第十一，寄存器分配是性能关键；寄存器过多降低 Occupancy，过少则可能引发 Spill。

第十二，volatile 主要约束编译器对内存访问的优化，不提供原子性或完整线程同步。

第十三，编译器会静态安排指令和进行软件流水化，以隐藏数据依赖和内存延迟。

第十四，优化可能增加指令或重复计算，以降低寄存器压力或提高整体吞吐率。

第十五，最终应结合 SASS、资源报告和实际 Profiling 判断优化是否有效，而不能只根据源代码或 Occupancy 下结论。

下一课将进入真正的 GPU 性能诊断：

第二十课：如何使用 Nsight Compute 分析一个 CUDA Kernel

重点包括：

Kernel Duration
SM Throughput
DRAM Throughput
Memory Workload Analysis
Scheduler Statistics
Eligible Warps
Warp Stall Reasons
Occupancy
Roofline Chart
如何判断计算受限、带宽受限或延迟受限
如何从性能报告定位到具体源码