简介：本文以 CUDA 矩阵乘法优化为主线，对比朴素实现与 2D 分块共享内存方案在访存效率、计算组织方式及适用场景上的差异。通过拆解 GPU 的 SIMT 架构、存储层级与分块策略，帮助读者理解矩阵乘法 kernel 的优化逻辑、访存瓶颈与选型依据，为后续对齐高性能计算库建立清晰的技术判断框架。
在大模型训练与推理中，矩阵乘法是计算密度最高、优化空间最大的算子之一。无论是全连接层、注意力机制中的 QK^T 计算，还是 Feed-Forward Network 中的权重投影，底层都依赖高效的矩阵乘实现。GPU 之所以能在这一场景中发挥巨大价值，核心在于其 SIMT 架构能够用大量线程并行执行同一份计算逻辑，同时通过多层存储体系缓解“计算快、访存慢”的矛盾。
然而，同样是矩阵乘法，不同的 kernel 实现方式在数据搬运、线程组织、共享内存利用和计算效率上差异极大。一个直观的对比是：朴素实现按“每个 block 计算一个输出元素”的方式组织，计算逻辑简单直接，但访存浪费严重；而2D 分块共享内存方案将输出矩阵切分为多个 tile，通过复用共享内存大幅减少对全局内存的重复访问，显著提升计算访存比。
本文围绕这两种实现方式进行对比分析，目的是帮助读者建立 CUDA 性能优化的核心思路：不是单纯“把循环写对”，而是要理解数据从全局内存到共享内存、再到寄存器、最终进入计算单元的流动路径，以及如何降低这条路径上的无效开销。
朴素实现采用最直接的线程映射方式：每个 block 负责计算输出矩阵 C 中的一个元素。block 内的多个线程分别计算 A 矩阵某一行与 B 矩阵某一列对应位置的部分乘积，再通过 block 内归约得到最终结果。
以公式 D = alpha · A · B^T + beta · C 为例，在 alpha=1、beta=0 的简化场景下，C 矩阵中每个元素 C[i][j] 等于 A 的第 i 行向量与 B 的第 j 行向量（因 B 被转置）的点积。朴素实现让每个 block 独立承担一次完整的点积计算。
这种方案的优点是逻辑清晰、索引计算简单，适合作为 CUDA 入门的第一版实现。但其访存模式存在明显问题：A 和 B 的同一块数据会被多个 block 重复从全局内存读取，而全局内存的带宽是 GPU 系统中相对稀缺的资源。
2D 分块方案改变了线程与数据的映射关系。它不再让每个 block 只计算一个输出元素，而是将输出矩阵 C 划分为若干个大小为 (b_m, b_n) 的 block tile。每个 block 负责计算一个完整 tile 的结果。同时，由于共享内存容量有限，无法一次性容纳整个 k 维度的数据，因此对 k 维度也进行切分，每个 block tile 内部沿 k 方向循环累加，每次处理 (b_m, b_k) × (b_k, b_n) 大小的子矩阵乘积。
这一方案的核心在于：先将 A 和 B 中将被重复使用的数据块从全局内存搬运到共享内存，后续计算直接从共享内存读取，避免反复访问 DRAM。共享内存的带宽远高于全局内存，且延迟更低。虽然共享内存的容量受限，但通过 k 维度的循环分块，可以让 kernel 适应任意尺寸的矩阵。
两种方案都运行在相同的 GPU 硬件之上，遵循 SIMT 的执行模型。所有线程执行同一份 kernel 代码，根据 blockIdx 和 threadIdx 确定自己在线程群中的位置，进而决定各自处理的数据范围。从这个意义上讲，两种方案的底层计算单元相同，都是利用 GPU 的大量 CUDA core 进行并行乘加运算。
在计算逻辑层面，两种方案最终都要完成相同数量的乘加操作：对于 m×n 的输出矩阵和 k 维的规约长度，总计算量为 m×n×k 次乘加。无论采用哪种 kernel 组织方式，计算本身的数学结果完全一致。
两者的数据流也共享同一基本路径：数据从全局内存（DRAM）出发，经过 L2 cache，进入 SM 内的共享内存或 L1 cache，最终到达寄存器供计算单元使用。理解这条路径，是优化任何一种 kernel 的前提。
访存是两种方案最本质的差异点。
朴素实现中，每个 block 独立完成一个输出元素的计算。以一个输出元素 C[i][j] 为例，它需要读取 A 的第 i 行（k 个元素）和 B^T 的第 j 列（也是 k 个元素）。如果有 m×n 个 block，每个 block 都从全局内存读取自己的那份 A、B 数据，那么整个计算过程的 DRAM 访问量为：
加载 A：m × k 行数据被 n 个 block 重复读取
加载 B^T：n × k 行数据被 m 个 block 重复读取
简化后，DRAM 访问总量约为 2 × m × n × k 次元素读取，加上 m × n 次输出写入。
2D 分块方案通过共享内存复用显著降低全局内存访问。在一个 block tile 内，大小为 (b_m, b_k) 的 A 子块和 (b_k, b_n) 的 B^T 子块被加载到共享内存后，可以被该 block 内所有线程反复使用。理想情况下，DRAM 访问量约为：
加载 A：m × k 次（每个元素只被对应 block tile 读取一次）
加载 B^T：n × k 次
写入 C：m × n 次
与朴素实现相比，分块方案的全局内存读取量从 2mnk 降低到 (m+n)k 级别。当 m、n 较大时，访存缩减效果非常显著。
朴素实现对共享内存的利用几乎为零。数据从全局内存读取后直接进入寄存器和计算单元，共享内存这一重要的片上存储资源被闲置。这导致计算单元经常处于等待数据到达的状态，因为全局内存的延迟远高于片上存储。
2D 分块方案则把共享内存作为数据复用的关键缓冲区。数据从 DRAM 搬运到共享内存的过程需要遵循合并访问原则，即同一个 warp 内的线程访问连续的 128 字节数据块，以保证最高的内存事务效率。数据进入共享内存后，计算阶段对共享内存的访问延迟远低于全局内存，使计算单元能够更充分地保持忙碌状态。
朴素实现的线程组织非常直观。每个 block 对应一个输出元素，block 内的线程通过归约完成向量点积。索引计算公式简单，代码可读性强，便于验证正确性。
2D 分块方案的线程组织则更为复杂。需要处理三层索引：block 级别的 tile 定位、block 内部线程到 tile 内元素的映射、以及 k 维度循环中的子块偏移。同时需要显式管理共享内存的加载、同步和复用。这种复杂度对应的是更高的性能上限，但在开发和调试上需要更多投入。
朴素实现中，block 内部需要一次归约来汇总部分乘积结果，使用共享内存或 warp shuffle 即可完成。这里存在一次 block 内的同步点。
2D 分块方案中，每个 k 维循环迭代都需要：先将 A、B 子块从全局内存加载到共享内存，执行 __syncthreads() 确保所有线程加载完成后才能开始计算；计算完当前子块后，还需要再次同步，确保所有线程读取完共享内存后才能被下一次加载覆盖。同步点的增加对数据依赖管理提出了更高要求，如果同步不当，容易产生数据竞争或读取到未更新数据的问题。
| 对比维度 | 朴素 matmul | 2D block tiling |
|---|---|---|
| 全局内存访问量 | 约 2mnk + mn | 约 (m+n)k + mn |
| 共享内存利用 | 几乎不使用 | 核心优化手段 |
| 数据复用方式 | 无跨 block 复用 | tile 内共享复用 |
| 线程映射复杂度 | 低 | 高 |
| 需要同步的次数 | 少（一次 block 内归约） | 多（每次 k 循环至少两次） |
| 对硬件带宽的依赖 | 强依赖全局内存带宽 | 降低全局内存带宽压力 |
| 代码可读性与调试难度 | 易读、易调试 | 结构复杂、调试成本高 |
| 性能上限 | 较低 | 接近高性能计算库的必要路径 |
当矩阵规模较小，或开发目标是验证 CUDA 编程基本概念、快速建立索引映射和线程模型的理解时，朴素实现是合理的起点。在小规模数据下，共享内存分块带来的收益不足以抵消代码复杂度的提升。此外，在原型验证阶段，首要任务是保证正确性，朴素实现能更快地验证算法逻辑。
当矩阵维度较大，计算访存比成为性能瓶颈时，分块共享内存方案是必要的优化步骤。尤其是大模型推理与训练中，矩阵乘法占据了大量计算时间，全局内存带宽往往成为首要制约因素。此时必须通过分块复用数据，将尽可能多的访存发生在片上共享内存中。
另外，对于需要在不同硬件上适配的 kernel，分块方案中 tile 大小的选取可以针对具体 GPU 的共享内存容量和计算单元数量进行调整，具备更好的可扩展性。
选择哪种实现方式，核心取决于两个因素：当前性能瓶颈的位置和开发团队对 CUDA 优化技术的掌握程度。
如果团队刚刚开始接触 GPU 编程，对 block、thread、warp 的组织方式还不够熟悉，建议先从朴素实现入手，完整走通“索引计算—数据读取—计算—写回”的流程。在此基础上，用性能分析工具观察全局内存访问量和计算单元占用率，再逐步引入共享内存分块。
如果团队已经具备基本的 CUDA 开发经验，并且明确目标是对齐高性能计算库的矩阵乘法性能，那么 2D 分块是绕不开的基础步骤。但这还不够，后续往往还需要进一步引入向量化访存、double buffering、warp tiling、register tiling 等技术，才能逐步逼近硬件理论峰值。
关键判断依据是：当全局内存带宽成为瓶颈、计算单元经常空闲等待数据时，优先考虑分块复用；当 kernel 的计算指令流本身成为瓶颈时，则需要从计算并行度、指令调度和寄存器使用等角度寻找优化空间。
从朴素实现迁移到 2D 分块方案，需要关注以下几个方面。
正确性验证。 分块方案涉及多层循环和多次同步，索引计算错误和同步遗漏都可能产生难以察觉的错误结果。建议先用小尺寸矩阵与 CPU 参考实现逐元素对比，确认无误后再扩大规模进行性能测试。
tile 大小的选择。 tile 大小直接影响共享内存占用和计算效率。tile 过小，无法充分复用数据；tile 过大，可能导致共享内存溢出，降低 SM 上可同时驻留的 block 数量，影响 occupancy。需要根据目标 GPU 的共享内存容量、寄存器资源和最大线程数进行实测调整。
合并访问的实现。 从全局内存向共享内存搬运数据时，要保证同一 warp 内的相邻线程访问相邻的内存地址。如果索引布局不当，合并访问退化为非合并访问，内存事务数会成倍增加，抵消分块带来的收益。
同步的正确使用。 每个 k 循环迭代中，加载共享内存后和重新写入共享内存前都需要恰当的同步。漏掉任何一个同步都可能导致数据竞争。同时，过度同步也会增加开销，需要在数据依赖正确的前提下合理控制同步点数量。
性能评估方法。 不同 GPU 架构、不同矩阵尺寸下，分块方案的最佳 tile 大小和循环组织方式可能不同。建议建立一套包含多种矩阵规模的基准测试集，在目标硬件上实测，避免仅凭理论分析做出结论。
朴素实现与 2D 分块共享内存方案代表了 CUDA 矩阵乘法优化路径上的两个重要阶段。前者解决了“如何用 GPU 并行计算矩阵乘法”的正确性问题，后者则回答了“如何让 GPU 的内存层次真正服务于计算复用”的性能问题。
两者的差异本质在于对访存效率的不同利用水平。朴素实现将大量时间花费在等待全局内存数据上，计算单元利用率低；2D 分块通过共享内存中的数据复用，显著降低了对全局内存带宽的依赖，提升了计算访存比，但也带来了线程组织复杂度和同步管理成本的上升。
对于 CUDA 开发者而言，理解这一对比的意义不仅在于掌握某一种 kernel 写法，更在于建立观察性能瓶颈的思维方式：先定位数据流动的瓶颈，再有针对性地调整存储层级利用策略。这也是后续继续深入优化矩阵乘法、并理解高性能计算库中更复杂优化技术的基础。