4、矩阵乘法基础：CPU矩阵乘法实现、GPU朴素矩阵乘法、性能瓶颈分析

好，咱们进入正题。矩阵乘法，说白了就是深度学习里最核心的计算单元。你想想看，全连接层、卷积层、注意力机制……底层拆开，全是一堆矩阵在乘来乘去。今天我就带你从最基础的CPU实现开始，一步步看到GPU上的朴素写法，再聊聊性能瓶颈到底卡在哪。

4.1 CPU上的矩阵乘法实现

先来个最直观的。假设我们有两个矩阵 A (M×K) 和 B (K×N)，要算 C = A × B。教科书上的定义就是三重循环：

// CPU上的朴素矩阵乘法
void matmul_cpu(float* A, float* B, float* C, 
                int M, int N, int K) {
    for (int i = 0; i < M; i++) {
        for (int j = 0; j < N; j++) {
            float sum = 0.0f;
            for (int k = 0; k < K; k++) {
                sum += A[i * K + k] * B[k * N + j];
            }
            C[i * N + j] = sum;
        }
    }
}

这段代码，我刚开始学CUDA时写了不下几十遍。逻辑很简单：C矩阵的每个元素，是A的一行和B的一列做点积。但注意，这里有个隐藏的坑——内存访问模式。

4.2 GPU朴素矩阵乘法：第一版尝试

好，现在把问题搬到GPU上。最直接的想法：每个线程算C矩阵的一个元素。线程块大小用16×16或者32×32，网格大小根据M和N来定。

// GPU朴素矩阵乘法（第一版）
__global__ void matmul_naive(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;
    }
}

这段代码，嗯，能跑。但性能嘛……我当年第一次跑的时候，拿一个1024×1024的矩阵测试，结果比CPU还慢。当时我就懵了，GPU不是并行神器吗？

4.3 性能瓶颈分析：为什么这么慢？

咱们来拆解一下。GPU的瓶颈主要来自三个方面：

4.3.1 全局内存访问延迟

每个线程都要从全局内存读A的一行和B的一列。全局内存的延迟是几百个时钟周期，而计算一个乘加只需要几个周期。说白了，计算单元大部分时间在等数据。

更糟糕的是，同一个线程块内的线程，访问的A和B地址是不连续的。比如线程(0,0)读A[0]和B[0]，线程(0,1)读A[0]和B[1]……这导致全局内存的访问无法合并（coalesced），带宽利用率极低。

4.3.2 重复读取数据

你想想看，A矩阵的同一行数据，会被N个不同的线程重复读取（每个列一个线程）。B矩阵的同一列数据，会被M个线程重复读取。这些数据从全局内存反复加载，完全没有复用。

以M=N=K=1024为例，每个线程要读K=1024次全局内存，总共1024×1024个线程，就是10亿次全局内存读取。而实际上，A矩阵只有1024×1024=100万个元素，B矩阵也是100万个。也就是说，每个数据被读了500次！

4.3.3 计算与访存比太低

对于每个输出元素，我们做了K次乘加（2K次浮点运算），但需要从全局内存读2K个数据。计算访存比大约是1（每读一个数做一次运算）。而GPU要发挥峰值性能，这个比值通常需要10以上。

指标	朴素GPU实现	理想值
全局内存访问次数	2 × M × N × K	M×K + K×N（仅一次）
计算访存比	1	≥ 10
内存合并效率	低（B矩阵按列访问）	高（连续访问）
数据复用	无	有（共享内存/寄存器）

4.4 小结：下一步怎么走？

看到这里，你应该明白了：朴素矩阵乘法的瓶颈不在计算，而在内存。GPU的计算单元其实很闲，它们都在等数据从全局内存慢悠悠地过来。

那怎么优化？我个人的经验是两条路：

• 利用共享内存：把A和B的块数据搬到片上共享内存，让一个线程块内的线程共享数据，减少全局内存访问。
• 优化内存合并：调整数据布局或访问模式，让同一warp的线程访问连续地址。

下一章，我会带你实现一个基于共享内存的分块矩阵乘法，性能直接翻几十倍。到时候你就知道，为什么说「共享内存是GPU优化的灵魂」了。