第四章：计算优化——ILP、Warp Level、Tensor Core与Matrix Core

各位同学，今天我们来聊聊计算优化。这个话题，说白了就是怎么让芯片算得更快。我这些年调过不少算子，踩过不少坑，今天把核心经验分享给你们。

4.1 指令级并行（ILP）——让流水线别闲着

指令级并行，听起来高大上，其实道理很简单：让GPU的流水线一直有活干。GPU的指令发射宽度有限，如果一条指令依赖上一条的结果，流水线就得等。

核心原则：减少数据依赖，让独立指令交错执行。

举个例子，我早期调一个向量加法的kernel，发现性能死活上不去。后来一看汇编，全是连续的load→add→store，流水线空转严重。改成这样就好了：

// 低效版本：连续依赖
float a = x[i];
float b = y[i];
float c = a + b;
z[i] = c;

// 高效版本：交错执行，隐藏延迟
float a0 = x[i];
float a1 = x[i+1];
float b0 = y[i];
float b1 = y[i+1];
float c0 = a0 + b0;
float c1 = a1 + b1;
z[i] = c0;
z[i+1] = c1;

你想想看，第二个版本里，load和add指令交错进行，流水线利用率高多了。我在项目中遇到过，ILP优化通常能带来10%-30%的性能提升，而且不花一分钱硬件成本。

小技巧：用__launch_bounds__控制线程数，给编译器更多ILP优化空间。我习惯把block size设成128或256，效果不错。

4.2 线程束级并行（Warp Level）——别让半个warp闲着

Warp是GPU执行的基本单位，32个线程一起跑。如果分支发散，部分线程就得等。嗯，这里要注意：warp内的分支越少越好。

我曾经调过一个图像滤波算子，里面有个if-else判断像素类型。结果性能惨不忍睹。后来改成这样：

// 糟糕：warp内分支发散
if (pixel_type == 0) {
    // 处理类型0
} else {
    // 处理类型1
}

// 优化：用select代替分支
float result = (pixel_type == 0) ? process_type0(data) : process_type1(data);

为什么？因为三元运算符会被编译成sel指令，所有线程都执行，只是选择结果。没有分支，warp就不会发散。

避坑指南：我曾经在warp内用__syncthreads()，结果死锁了。记住，__syncthreads()是block级别的同步，warp内不需要它——因为warp本身就是同步执行的。

还有一个技巧：warp shuffle指令。用__shfl_xor_sync、__shfl_down_sync这些，可以在warp内直接交换数据，不用走shared memory。我调reduce算子时，用shuffle比用shared memory快了30%。

// warp内归约，用shuffle
float val = input[threadIdx.x];
for (int offset = 16; offset > 0; offset >>= 1) {
    val += __shfl_xor_sync(0xffffffff, val, offset);
}
if (threadIdx.x == 0) output[blockIdx.x] = val;

4.3 Tensor Core与Matrix Core——专用硬件的威力

Tensor Core（NVIDIA）和Matrix Core（AMD）是专门为矩阵乘法设计的硬件单元。它们不是通用ALU，而是一次算一个4x4或16x16的小矩阵。

我刚开始用Tensor Core时，以为直接调cuBLAS就行。后来发现，手动调wmma指令能榨出更多性能。比如：

// CUDA Tensor Core示例（wmma）
nvcuda::wmma::fragment<nvcuda::wmma::matrix_a, 16, 16, 16, half, nvcuda::wmma::row_major> a_frag;
nvcuda::wmma::fragment<nvcuda::wmma::matrix_b, 16, 16, 16, half, nvcuda::wmma::col_major> b_frag;
nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> c_frag;

nvcuda::wmma::load_matrix_sync(a_frag, a, 16);
nvcuda::wmma::load_matrix_sync(b_frag, b, 16);
nvcuda::wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
nvcuda::wmma::store_matrix_sync(c, c_frag, 16, nvcuda::wmma::mem_row_major);

AMD的Matrix Core类似，用__builtin_amdgcn_mfma_f32_16x16x16f16这类内建函数。我个人习惯：能用Tensor Core就别用普通CUDA Core，能快5-10倍。

关键点：Tensor Core要求数据布局对齐（通常是16x16的tile），而且精度有限（FP16、BF16、INT8等）。如果你的模型精度要求高，得做混合精度训练。

4.4 三种并行的协同使用

实际项目中，这三种并行不是孤立的。我调过一个卷积算子，是这样组合的：

ILP层面：循环展开4次，隐藏load延迟
Warp层面：用shuffle做channel维度的归约
Tensor Core层面：用wmma做4x4的矩阵乘法

结果比原始版本快了8倍。你想想看，优化不是单点突破，而是系统性的。

我的经验：先用profiler看瓶颈在哪。如果是计算瓶颈，优先上Tensor Core；如果是访存瓶颈，先优化ILP和warp调度。别一上来就全上，容易翻车。

4.5 知识体系总览

下面这张图，是我自己总结的优化框架。你可以把它当成checklist：

这张图把三种并行方式的关系讲清楚了。我个人习惯把它贴在工位上，每次调优前先看一眼，提醒自己别漏了哪个维度。

4.6 实战建议

最后，给几个实操建议：

先测基线：用nvprof或rocprof跑一遍，看occupancy、IPC、带宽利用率
ILP优先：如果IPC低于理论值一半，先做ILP优化
Warp次之：检查分支发散，用__ballot_sync诊断
Tensor Core压轴：如果计算是瓶颈，且精度允许，果断上Tensor Core

注意：Tensor Core不是万能的。我见过有人在小矩阵上硬用wmma，结果因为数据搬运开销太大，反而比普通CUDA Core慢。记住，矩阵尺寸至少16x16才有收益。

好了，这一章就到这里。记住，优化是门手艺活，多调多试，慢慢就有感觉了。

想要更多资料，联系可可

📚 课程目录

第四章：计算优化——ILP、Warp Level、Tensor Core与Matrix Core

4.1 指令级并行（ILP）——让流水线别闲着

4.2 线程束级并行（Warp Level）——别让半个warp闲着

4.3 Tensor Core与Matrix Core——专用硬件的威力

4.4 三种并行的协同使用

4.5 知识体系总览

4.6 实战建议