并行计算模式：从向量加法到矩阵乘法

各位同学，今天我们来聊聊并行计算模式。说实话，这是CUDA编程里最基础也最核心的内容。我刚开始接触GPU编程时，总觉得并行计算很玄乎，后来发现，说白了就是把一个大任务拆成很多小任务，让GPU里成千上万个线程同时干活。

向量加法：Hello World级别的并行

向量加法是并行计算的入门案例。两个数组对应元素相加，这个操作天然就是并行的——每个元素的计算互不依赖。

__global__ void vecAdd(float *A, float *B, float *C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

这段代码看起来简单，但有个细节我吃过亏。你看那个if (i < N)的判断，为什么要加？因为数据总量不一定能被线程块大小整除。我曾经在项目里忘了加这个边界检查，结果数组末尾多写了一些垃圾数据，排查了半天。

调用方式是这样的：

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里有个小技巧：用(N + threadsPerBlock - 1) / threadsPerBlock来计算线程块数量，可以保证线程总数不少于数据量，配合边界检查就万无一失了。

矩阵乘法：并行计算的经典案例

矩阵乘法比向量加法复杂得多。每个输出元素是A的一行和B的一列的点积。我当年面试时，面试官让我手写CUDA矩阵乘法，写出来不难，但要写出高性能版本，那才是真功夫。

朴素实现：每个线程计算一个输出元素。

__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;
    }
}

为什么会这样？因为全局内存的延迟很高，大约几百个时钟周期。而计算一个点积需要K次乘加操作，每次都要等数据从全局内存过来，GPU的计算单元大部分时间都在空转。

优化思路：使用共享内存，把数据分块加载到片上缓存。

__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 row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    
    float sum = 0.0f;
    for (int tile = 0; tile < (K + TILE_SIZE - 1) / TILE_SIZE; tile++) {
        // 协作加载数据到共享内存
        if (row < M && tile * TILE_SIZE + threadIdx.x < K)
            As[threadIdx.y][threadIdx.x] = A[row * K + tile * TILE_SIZE + threadIdx.x];
        else
            As[threadIdx.y][threadIdx.x] = 0.0f;
            
        if (col < N && tile * TILE_SIZE + threadIdx.y < K)
            Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
        else
            Bs[threadIdx.y][threadIdx.x] = 0.0f;
        
        __syncthreads();
        
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        
        __syncthreads();
    }
    
    if (row < M && col < N) {
        C[row * N + col] = sum;
    }
}

性能分析工具：nvprof与Nsight Systems

写完了代码，怎么知道它跑得快不快？靠猜可不行。我见过太多人凭感觉优化，结果越优化越慢。这时候就需要性能分析工具了。

nvprof：命令行利器

nvprof是NVIDIA提供的命令行性能分析工具。用法很简单：

nvprof ./my_cuda_program

它会输出每个kernel的执行时间、内存传输时间、占用率等信息。我个人习惯先跑一遍nvprof，看看整体情况。

更详细的用法：

nvprof --metrics gld_efficiency,gst_efficiency,occupancy ./my_cuda_program

这个命令会告诉你全局内存加载/存储效率，以及SM的占用率。我曾经优化一个矩阵乘法，发现gld_efficiency只有30%多，说明大部分访存带宽被浪费了。后来改用共享内存分块，效率提升到了80%以上。

指标	含义	理想值
gld_efficiency	全局内存加载效率	接近100%
gst_efficiency	全局内存存储效率	接近100%
occupancy	SM占用率	50%以上
achieved_occupancy	实际达到的占用率	越高越好

Nsight Systems：可视化分析

Nsight Systems是图形化的性能分析工具，比nvprof更直观。它能展示时间轴上的kernel执行、内存拷贝、CPU活动等。

启动方式：

nsys profile -o output_file ./my_cuda_program

然后打开生成的.nsys-rep文件，就能看到漂亮的时序图了。

知识体系结构图

下面这张图展示了本章的核心知识脉络：

总结

向量加法是入门，矩阵乘法是进阶，性能分析工具是武器。这三者构成了CUDA并行编程的基础框架。我个人建议的学习路径是：先跑通向量加法，理解线程模型；再挑战矩阵乘法，体会共享内存的威力；最后用nvprof和Nsight Systems来验证你的优化是否有效。

记住一句话：没有性能分析就没有优化。不要凭感觉调优，让数据告诉你瓶颈在哪里。

想要更多资料，联系可可