4、CUDA并行计算基础：向量加法实现、矩阵乘法实现、性能分析工具nvprof/Nsight Systems初探
好，咱们直接进入正题。这一章是CUDA并行计算的“手把手”环节。说白了，就是让你亲手写出第一个能跑在GPU上的程序，然后看看它到底有多快。
我个人习惯，学任何并行框架，先搞定两个例子：向量加法和矩阵乘法。前者让你理解“线程怎么干活”，后者让你理解“线程怎么协作干活”。搞懂了这两个，后面再学卷积、Transformer什么的，心里就有底了。
4.1 向量加法：Hello World的GPU版
先来个最简单的。两个数组对应元素相加。CPU上写个for循环就完事，GPU上怎么写？
嗯，这里要注意：GPU是“单指令多线程”的架构。你写一个函数（叫kernel），然后告诉GPU：“给我启动N个线程，每个线程执行这个函数”。每个线程通过内置变量 threadIdx.x 知道自己是第几个。
核心思想： 一个线程处理一个元素。N个元素，就启动N个线程。
// 向量加法的kernel函数
__global__ void vecAdd(float *A, float *B, float *C, int N) {
int i = threadIdx.x + blockIdx.x * blockDim.x;
if (i < N) {
C[i] = A[i] + B[i];
}
}
// 调用方式（假设N=1024，每个block 256个线程）
int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);你看，代码就这么几行。但坑在哪？我刚开始学的时候，老忘记做边界检查。万一N不是blockSize的整数倍呢？最后几个线程就会访问越界。所以那个 if (i < N) 是保命符。
我的经验： 写kernel时，第一件事就是算好全局线程ID，第二件事就是做边界检查。养成肌肉记忆，能省掉你后面无数debug时间。
4.2 矩阵乘法：从单线程到共享内存
向量加法太简单了，体现不出GPU的威力。矩阵乘法才是重头戏。你想想看，一个1024x1024的矩阵相乘，CPU上要算10亿次乘加。GPU怎么加速？
最朴素的思路：每个线程算输出矩阵中的一个元素。也就是一个线程负责一行和一列的点积。
__global__ void matMulNaive(float *A, float *B, float *C, int M, int N, int K) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < M && col < N) {
float sum = 0.0f;
for (int k = 0; k < K; k++) {
sum += A[row * K + k] * B[k * N + col];
}
C[row * N + col] = sum;
}
}这个版本能跑，但性能很差。为什么？因为每个线程都在反复从全局内存读数据。A矩阵的一行会被读N次，B矩阵的一列会被读M次。全局内存的带宽是有限的，这成了瓶颈。
我在项目中遇到过类似问题。当时算一个1280x1280的矩阵乘法，用这个朴素版本跑了快2秒。后来优化到用共享内存，直接降到80毫秒。差距就是这么大。
避坑指南： 我曾经天真地以为GPU算得快，随便写写就能起飞。结果发现，如果不考虑访存模式，性能可能还不如CPU。记住：GPU的算力远大于内存带宽，所以优化访存永远是第一位的。
优化的关键：用共享内存做数据复用。把A的一块子矩阵和B的一块子矩阵加载到共享内存里，然后让一个block内的线程共享这些数据。
// 使用共享内存的矩阵乘法（分块版本）
#define TILE_SIZE 16
__global__ void matMulShared(float *A, float *B, float *C, int M, int N, int K) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int bx = blockIdx.x, by = blockIdx.y;
int tx = threadIdx.x, ty = threadIdx.y;
int row = by * TILE_SIZE + ty;
int col = bx * TILE_SIZE + tx;
float sum = 0.0f;
for (int tile = 0; tile < (K + TILE_SIZE - 1) / TILE_SIZE; tile++) {
// 协作加载数据到共享内存
if (row < M && tile * TILE_SIZE + tx < K)
As[ty][tx] = A[row * K + tile * TILE_SIZE + tx];
else
As[ty][tx] = 0.0f;
if (col < N && tile * TILE_SIZE + ty < K)
Bs[ty][tx] = B[(tile * TILE_SIZE + ty) * N + col];
else
Bs[ty][tx] = 0.0f;
__syncthreads(); // 确保所有线程都加载完毕
for (int k = 0; k < TILE_SIZE; k++)
sum += As[ty][k] * Bs[k][tx];
__syncthreads(); // 确保所有线程都算完，再加载下一块
}
if (row < M && col < N)
C[row * N + col] = sum;
}这里有个关键点：__syncthreads()。它是个同步屏障。我刚开始写的时候，经常忘记加这个，结果数据还没加载完就开始算，算出来的全是错的。嗯，这个坑我踩过不止一次。
4.3 性能分析工具：nvprof与Nsight Systems
代码写完了，跑起来了，结果也对。然后呢？你得看看它到底跑得多快，瓶颈在哪。这就轮到性能分析工具上场了。
老牌工具是 nvprof，新工具是 Nsight Systems。我个人的建议：新项目直接用Nsight Systems，nvprof已经逐渐被弃用了。但有些老项目还在用，所以两个都得会。
4.3.1 nvprof快速上手
nvprof用起来很简单，直接在命令行里包一下你的程序就行：
nvprof ./my_cuda_program它会输出每个kernel的耗时、内存传输时间、占用率等信息。我最常用的几个参数：
--print-gpu-trace：显示每个GPU操作的详细时间线--metrics：查看具体指标，比如--metrics gld_efficiency看全局内存加载效率
关键指标： 全局内存加载效率（gld_efficiency）。如果低于80%，说明你的访存模式有问题，需要调整数据布局或使用共享内存。
4.3.2 Nsight Systems：可视化分析
Nsight Systems是个图形化工具，能看到整个程序的时间线。CPU在干什么，GPU在干什么，数据传输和kernel执行是不是重叠的，一目了然。
我一般这样用：
- 先跑一次
nsys profile -o output_file ./my_program生成分析文件 - 然后用
nsys-ui打开，看时间线 - 重点关注：GPU利用率、内存拷贝时间、kernel启动开销
记得有一次，我发现我的程序GPU利用率只有30%。一看时间线，原来是CPU在忙着准备数据，GPU一直在空等。后来改成异步传输和流水线，利用率直接拉到90%以上。
小技巧： 用Nsight Systems看时间线时，注意看CPU和GPU之间有没有“空隙”。如果有，说明你的程序没有做到CPU和GPU的并行工作。用CUDA Stream可以解决这个问题。
4.4 本章知识体系
下面这张图，是我自己梳理的本章核心逻辑。从最简单的向量加法，到需要协作的矩阵乘法，再到性能分析工具，是一条完整的链路。
说白了，这一章就是让你把CUDA的“三板斧”练熟：怎么写kernel、怎么优化访存、怎么用工具找瓶颈。这三样东西，后面每一章都会用到。
好，这一章就到这。代码示例我都跑过，你直接复制到项目里就能用。如果遇到问题，先检查边界条件，再检查同步，最后用工具看瓶颈。这三步走下来，90%的问题都能解决。