企业官网建设流程全解析

做高性能矩阵计算的人，应该都遇到过同一个场景：明明把计算量算得很清楚，但程序一跑起来，瓶颈根本不在浮点运算单元，而在内存搬运。尤其是大矩阵场景，数据搬进搬出的时间远超计算本身，GPU 一堆 SM 空转等着数据到位。我最近一年一直在围绕这个问题做优化，最后把思路收敛到一个叫 KDA（Kernel-based Decomposition Algorithm，核分解算法）的方案上。这个方案把矩阵记忆策略和并行计算调度串成一条完整的链路：先规划矩阵数据在内存层次中的存放位置，再按核把矩阵运算拆解成大量可并行的子任务，最后通过依赖图做负载均衡调度。这篇文章我会把 KDA 从原理到落地讲清楚，包括矩阵记忆的具体设计、核分解的数学拆解、任务图调度方式，以及我实际跑过的基准测试和踩坑记录。如果你也在做大矩阵运算或并行计算优化，里面有不少可以直接抄走的思路。

1. 为什么矩阵访问模式决定了并行计算的上限

1.1 从 Roofline 模型看计算与内存的关系

搞并行优化的人嘴里常挂着一个词叫“算术强度”，用大白话说就是“每搬一个字节的数据，能换来多少次浮点运算”。这个比值决定了一个计算任务是计算密集型还是访存密集型。拿最经典的 GEMM（矩阵乘法）来说，假设矩阵规模是N×N，总计算量大约2N³次浮点运算。如果把A、B、C三个矩阵一次性放进内存，数据搬运量大约是3N²字节。算一下算术强度：

(2N³) / (3N²) = 2N / 3

当N=2048时，这个值大约在 1365 FLOP/Byte 左右，听起来非常“计算密集”。但前提是数据能一直被留在缓存或共享内存里复用。如果代码写得不行，每算一个输出元素都把源数据从主存重新拉一遍，那实际算术强度会断崖式下跌，搞不好连 5 FLOP/Byte 都不到。

Roofline 模型里画得很清楚：硬件有一个计算峰值，也有一个内存带宽上限，两者相除得到一个“拐点”。算术强度高于拐点，程序受计算峰值限制；低于拐点，程序就死死卡在带宽上。很多大矩阵程序跑不快，根本不是 CPU 或 GPU 不够强，而是访问模式把算术强度压到了拐点以下，算力再高也白搭。

1.2 大矩阵的三个层次瓶颈

大矩阵到底贵在哪？我习惯把它拆成三层来看：

第一层是容量。一个8192×8192的float32矩阵，裸大小是 256MB。如果还有 B、C、中间结果，轻松上 GB 量级。而 GPU 的共享内存一个 SM 通常只有几十到两百多 KB，片上缓存也有限。数据放不下，就只能被反复换进换出。

第二层是带宽。主存带宽看着很高，但多核同时抢的时候，每个核分到的实际带宽会明显下降。更麻烦的是，如果访问模式是“跳着读”，比如列主序矩阵按行访问，每次只取一个 float，那一次 32 字节甚至 64 字节的缓存行大部分都被浪费了。实际有效带宽可能只有理论值的 1/4。

第三层是延迟。DRAM 的一次访问延迟在几十到上百纳秒量级，计算单元执行一条浮点指令只要几个周期。一旦数据没有提前到位，ALU 就只能空等。现代架构要靠大量并行掩盖延迟，但掩盖的前提是访存请求能被流水化地发出去，而流水化依赖访问模式的规律性。

这三层问题，本质上都指向同一件事：矩阵数据放在哪里、以什么顺序被访问，直接决定了并行的真实效率。这也是“矩阵记忆”这个思路的出发点。

1.3 经典矩阵分解为什么没能解决并行访存问题

说到矩阵分解，很多人第一反应是 LU、SVD、QR 这些经典数值算法。它们解决的是“怎样把一个矩阵变成更易用的结构”，比如把方程组求解变成三角化求解，或者把矩阵压缩成奇异值形式。这些分解在数值代数上很漂亮，但对并行访存的帮助有限。

原因很简单：经典分解的目标是改变矩阵的数学结构，而不是改变数据在存储层次中的搬运方式。SVD 可以把一个8192×8192的矩阵近似成低秩形式，内存确实省了，但每一个 SVD 步骤内部仍然是密集矩阵运算，照样面临访存瓶颈。LU 分解里有大量行交换和秩更新，切块之后通信模式反而更复杂。

我并不是说经典分解没用，而是说在并行计算这个维度上，它们和“为并行访存而设计的分解”根本不是一回事。KDA 从一开始就把目标定在存储层次上：不是改变矩阵的数值性质，而是改变矩阵被切分、放置、复用和调度的方式。

2. 矩阵记忆：决定数据该住在哪一层

2.1 矩阵记忆不是简单的“存下来”，而是布局编排

“矩阵记忆”这个词，我最初理解得特别浅，以为就是把大矩阵拆成小块放进缓存。后来在实践中才意识到，矩阵记忆真正要做的是回答三个问题：

• 数据应该以什么顺序存在于主存中？
• 哪些数据应该被搬进片上存储并长期驻留？
• 哪些数据只需要流式读取一次？

这三个问题如果答对了，并行计算的开销能降一个数量级。我之前遇到过一段很典型的代码：做矩阵乘法前先对 B 做转置，代码逻辑完全正确，但 B 转置后仍然是按列访问的存储布局，GPU 上每个 warp 读数据时跨步访问，缓存命中率惨不忍睹。后来改成“B 也按分块转置存储，并且每块连续排列”，整体耗时直接掉了 30% 以上。布局的差异就是这么明显。

在 KDA 的语境里，矩阵记忆指的是：根据运算模式和硬件存储层次，主动决定矩阵的物理排布形式。它包含三个层面的决策：内存排列（行主序还是列主序，还是分块排列）、缓存驻留（哪些分块高频复用，需要固定在共享内存或 L2 中）、冗余策略（边界数据是否需要镜像复制以降低通信）。

2.2 矩阵记忆的三种基本实现手段

我用得最多的三种手段是重排、分块和稀疏化。

重排最典型的就是转置布局。矩阵 B 在C = A × B里是按“列”访问的，如果原始数据是行主序，那就把它预先转置成列主序，或者干脆以分块转置后的格式存储。这样每次读取的是一段连续地址。

分块就是把大矩阵切成小的 tile。这里的关键不是切多大，而是切完之后的放置策略。高频参与内层循环的 tile 要尽量放在高速存储中；低频数据则留在主存里，用预取指令提前拉。我在实验里常用的分块大小是 64×64，由后文共享内存尺寸计算会说明原因。

稀疏化则是针对“大量结构性重复”的矩阵，比如对角占优矩阵、带形矩阵或图邻接矩阵。这类矩阵直接按稠密格式存，浪费严重；用稀疏格式存，再配合分块索引表，能在不改变分解精度的前提下显著压缩访存体量。

三类手段可以叠加使用。一套完整的矩阵记忆方案，最终会输出一个“布局计划（Layout Plan）”，告诉后续调度层每个数据块住在哪、什么时候搬、搬多少。

2.3 一个便于对比的存储账本

我习惯用一张表来估算不同存储方案的实际开销。以单精度矩阵为例：

这里 KDA 的“分块+驻留计划”不是做数据压缩，而是通过只保留“参与计算的分块”和“跨块边界需要的镜像数据”，避免把整个矩阵无差别地塞进显存。边界镜像的开销和分块数量有关，块越小镜像比例越高，所以分块大小不能无限缩小。

3. KDA 核心原理：按核分解的三个阶段

3.1 什么是“核”

KDA 里的“核”不是 CUDA kernel 那种完整函数，而是一个更小的计算单元：一组输入数据、一组输出数据、以及作用在其上的运算规则。可以把它理解成“数据块 + 算子”的绑定关系。一个矩阵乘法可以拆成几百个核，每个核只负责一小块输出的计算。

核的粒度必须大于单线程，否则调度开销太大；也必须小于整个矩阵，否则无法并行。我在多数场景选择 64×64 的核粒度，原因后面会细说。

3.2 第一阶段：记忆规划

先把整个运算拆成数据块的访问图。以C = A × B为例，把 A 按行分成M/b块，B 按列分成N/b块，K 维按深度分成K/b段。于是：

C(i,j) = Σ_k A(i,k) × B(k,j)

这里有(M/b)×(N/b)×(K/b)次小块的乘法组合。记忆规划阶段要统计每个小块被访问的频次。很明显，C 的每个输出块会被累计K/b次；A 的每个行块会被N/b个输出块重复读取；B 的每个列块同理。

访问频次越高，越值得把数据“记忆”在片上存储里。频次低的，则用流式加载。KDA 在这一步输出的是一张“驻留表”：哪些数据块放入共享内存或 L2 缓存，哪些块只做一次性读取。

3.3 第二阶段：核生成

有了驻留表，下一步就是把运算拆成核。一个核的定义包括三个信息：输入数据块的引用、算子类型、输出数据块。比如一个核可以写为：

K(bx, by, k) : C(bx, by) += A(bx, k) × B(k, by)

KDA 的核生成不是简单平均切块，而是根据驻留表自适应调整核的边界。高频区域可以把核适当做大，让更多复用发生在片上；低频区域核做小，减少无谓的片上占用。换言之，KDA 的核网格可以是“非均匀”的，这也是它和传统等宽分块（Tiling）最大的区别之一。

以 4×4 输出块为例，假设 K 维只有 2 段，那么总共生成4×4×2 = 32个核，但其中 16 个 C 输出块各自被两个核顺序更新，因此核之间存在“同一输出块的先后依赖”。这个依赖关系被递交给第三阶段。

3.4 第三阶段：构建依赖图

依赖图是一个有向无环图（DAG），节点是核，边是数据依赖。在我的场景里主要关心三种依赖：

• RAW（读后写）：后一个核要读取前一个核刚写出的数据。比如同一个 C 块在 k=0 被更新后，k=1 的核必须等它完成。
• WAR（写后读）：后一个核要覆盖数据，而前一个核还在读。这种依赖通常通过分配独立临时空间来消除。
• WAW（写后写）：两个核写同一位置。需要对写操作排序或用原子操作。

其中 RAW 是真正影响调度的；WAR 和 WAW 大多数时候可以在矩阵记忆规划时通过“每块独立累加”来规避——给每个核一个独立的临时输出缓冲区，最后再做一次归并。这样做牺牲一点存储，换来了调度自由度，我觉得非常值得。

4. 从依赖图到并行调度：把核塞满每一颗计算单元

4.1 调度器先处理依赖，再谈性能

依赖图建好之后，理想调度是：尽可能让没有依赖关系的核同时执行，同时让有依赖关系的核保持足够的顺序。对于 GEMM 来说，依赖关系主要沿着 K 维。只要调度器保证同一个 C 块的 k 段按顺序执行，不同 C 块之间完全可以乱序并行。

一个非常实用的做法是“按 C 块分组，按 k 段排序”。调度器做两轮：第一轮按 C 块坐标生成批次，第二轮在批次内按 k 索引递增调度。这样既能保证局部性，又能让所有 C 块在宏观上齐头并进，避免某些计算单元忙死、某些闲死。

4.2 静态调度 vs 动态调度

静态调度就是提前把所有核分配给固定的计算单元，好处是零运行时开销，坏处是负载可能不均。矩阵数据如果分布不均匀（比如稀疏矩阵里有些块非零元素多、有些块几乎全零），固定的静态映射会产生明显短板。

动态调度则相反，每个计算单元干完活就去队列里取下一个核，相当于“工作窃取”。在 GPU 上我通常用原子计数器维护一个待执行队列：每个 block 启动时从计数器里取一段核索引。开销大概每个核几次原子操作，但负载均衡效果显著。

我自己的经验是：稠密矩阵、结构规整的用静态调度；稀疏矩阵、非零结构不规则的用动态调度。KDA 的依赖图本身不变，变的只是“谁来执行核”的策略。

4.3 跨核边界的通信怎么控制

分块之后，边界数据不可避免会被相邻核重复读取。以 64×64 的核为例，如果做的是纯 GEMM，A 的行块和 B 的列块在每个核的二维网格里会被重复拉取，这部分属于“跨核复用”。处理办法是把跨核复用的数据在矩阵记忆阶段就复制到共享存储里，让相邻核从片上读取而不是重新访问主存。

另一个重点是把通信集中在长边而不是短边。核越大，通信/计算比越低。一个b×b的核，边界开销是O(b)，计算量是O(b³)，所以增大 b 能显著摊薄通信成本。这也是为什么核粒度不能太小——太小虽然并行度高，但通信开销会把你省下的时间又吃掉。

5. 基于 CUDA 的 KDA 落地实现

5.1 为什么我选了 CUDA 而不是 MPI 或 OpenMP

我做 KDA 的主要场景是大矩阵的稠密和半稠密运算，这类负载天然适合 GPU 的 SIMT 模型。MPI 适合跨节点分布式，但对单卡内部的共享存储层次控制不够细；OpenMP 在 CPU 上表现好，但对 GPU 的共享内存、寄存器级的调度控制力不足。CUDA 的优势是能让我精确控制共享内存布局和 block 调度，跟 KDA 的矩阵记忆理念非常匹配。

当然如果后面要扩展到多 GPU 或多节点，KDA 的依赖图本身是可序列化的，完全可以作为 MPI 的上层调度框架。但那部分我还没彻底做完，暂时先聚焦单卡。

5.2 核心数据结构和伪代码

KDA 的工程实现我分成三个模块：矩阵记忆规划器、核依赖图构建器、执行调度器。核心数据结构大概是这样的：

struct KernelPlan { int block; // 核边长 int mi, mj, kparts; // 核网格的三个维度 Layout layout; // RowMajor / ColumnMajor / Tiled vector<int> access_freq; // 每个数据块的预期访问频次 vector<int> resident_flags; // 标记哪些块常驻高速存储 }; struct KernelNode { int bx, by, k; vector<int> deps; // 依赖的核节点索引 }; struct KernelGraph { vector<KernelNode> nodes; vector<vector<int>> adj; // 邻接表 };

矩阵记忆规划器的伪代码可以简化为：

void build_kda_plan(const Matrix& A, const Matrix& B, Matrix& C, int block, KernelPlan* plan) { plan->block = block; plan->mi = (A.rows + block - 1) / block; plan->mj = (C.cols + block - 1) / block; plan->kparts = (A.cols + block - 1) / block; // 统计每个输出块的访问频次 plan->access_freq.assign(plan->mi * plan->mj, 0); for (int i = 0; i < plan->mi; i++) for (int j = 0; j < plan->mj; j++) plan->access_freq[i * plan->mj + j] = plan->kparts; // 高频块标记为常驻 plan->resident_flags.assign(plan->mi * plan->mj, 0); for (int idx = 0; idx < plan->mi * plan->mj; idx++) if (plan->access_freq[idx] > kparts / 2) plan->resident_flags[idx] = 1; }

别小看这个结构，依赖图的构建和调度器都以这个 plan 为基础。是否驻留、核大小、访问频次，全在计划阶段定好，运行时不再做任何动态决策。

5.3 共享内存的尺寸计算与指针划分

前面反复说 64×64 是我常用的核粒度，原因在共享内存这里一目了然。CUDA 的一个线程块通常能用的共享内存上限是 48KB（较新架构可以申请更多，但要额外配置）。如果每个核需要存放 A 块、B 块和 C 块各一份：

64 × 64 × 4 字节 × 3 = 48KB

刚好卡在绝大多数 GPU 的默认共享内存上限上，一分不浪费。如果块大小改成 128，那单算 A 块就是 64KB，直接超标，必须把块切小或申请动态共享内存。而我实测发现，64 的粒度在占用率和数据复用之间平衡得非常好。

分配时我把共享内存声明成动态共享内存，用指针分区：

extern __shared__ float smem[]; // 内存布局：前 16KB 给 A 块，接着 16KB 给 B 块，最后 16KB 给 C 块 float* tileA = smem; float* tileB = smem + block * block; float* tileC = smem + 2 * block * block;

注意动态共享内存的起始地址是按 16 字节对齐的，所以三个 16KB 区域的起始地址天然对齐到 4 字节以上，float 访问没问题。如果将来要换 double，建议强制按__align__(16)做一次对齐调整。

5.4 容易被忽略的 B 矩阵转置

实现 GEMM 时，新手最容易忽略的就是 B 矩阵的访问模式。C 的每个输出元素对应到 B 的一整列，而列访问在行主序存储下是跨步的。GPU 的缓存行一般是 128 字节，跨步读一列意味着每 128 字节里可能只有一个 float 被用到，剩下全浪费。

KDA 的矩阵记忆规划里，我会把 B 预转置成“Tiled 存储”：按 64×64 分块，块内转置，块与块连续存放。这样同一个核读 B 块时，地址是连续推进的，warp 内 32 个线程读同一块数据的效率直接拉满。代价是预处理时多一次转置的开销，但对于大数据量的多次反复计算，这笔开销摊到几千次核执行里，基本可以忽略。

6. 实测数据与踩坑记录

6.1 测试环境和参数

我在自己常用的 GPU 节点上跑了三组测试，环境是老一代的 V100 和一块消费级 RTX 3080，做对照。矩阵规模选4096×4096和8192×8192，数据类型 float32。对照组分别是：朴素三重循环实现、经典 64×64 等宽分块实现、KDA 实现，以及同一台机器上的 cuBLAS 作为参考上限。

以下数据只代表我这台机器的结果，硬件不同会变化，但相对趋势是稳定的：

KDA 没能超过 cuBLAS，这在我的预期内，毕竟 cuBLAS 用了大量手工汇编级优化和更细的调度策略。但 KDA 在纯 GEMM 场景下能跑到 cuBLAS 的 83% 到 86%，同时还能处理 cuBLAS 不太擅长的复杂算子组合和稀疏结构，我觉得这个性价比已经很高了。

6.2 踩坑一：共享内存 48KB 上限导致占用率崩了

最早的版本我把核块设为 128×128，想减少块数量、提高复用效率。结果共享内存需求是128×128×4×3 = 192KB，远超默认限制。后来通过cudaFuncAttributeMaxDynamicSharedMemorySize申请了 200KB，性能不但没提升，占用率反而掉到惨不忍睹：一个 SM 只能驻留一个线程块，分块内部的并行度全被调度器给压没了。

改成 64×64 之后，共享内存占用 48KB，一个 SM 能驻留多个块，占用率回到正常水平，延迟掩盖能力明显恢复。这个教训就是：核越大复用效率越高，但共享内存一旦顶爆，整个并行度会断崖式下跌，最终得不偿失。

6.3 踩坑二：静态调度在稀疏矩阵上出现长尾效应

我在处理一个带形矩阵时，静态调度的 KDA 耗时突然增加 40%。排查后发现，某些 C 块对应的非零段非常少，核执行速度快；另一些块几乎全程参与，核执行速度慢。静态调度把所有核均匀分到计算单元上，结果慢的核拖了整个批次的尾巴，大部分计算单元在干等。

解决办法就是动态调度：每个 block 从全局任务队列里取下一个核组，快的块自然多干几个，慢的块少干几个。替换后就这一招，耗时降了 33%。现在我的默认策略是：稠密矩阵静态调度，稀疏或半稀疏矩阵动态调度。

6.4 踩坑三：任务图顺序没考虑数据亲和性

有一版调度器严格按依赖图拓扑排序发派任务，没有任何逻辑错误，可实测缓存命中率一直不高。后来细看才发现，拓扑序只保证了“依赖正确”，没保证“空间相邻的核尽量同时执行”。两个计算相邻数据块的核被派到了不同批次，中间隔了几十万个时钟周期，缓存早被别的数据冲掉了。

我在调度器里加了一个二级策略：先在空间上按 C 块坐标分组，同一大块内的核尽量相邻调度；再在组内部用拓扑序排序。这个策略执行后，L2 命中率提升明显，整体耗时又降了 15%。现在想想，这个道理其实和 CPU 上的 NUMA 调度很像——逻辑上的依赖只是底线，物理上的位置才是性能。

7. KDA 的适用边界，以及我后续打算怎么做

KDA 适合什么场景？我目前的判断是：矩阵规模足够大、有明确可复用的数据块、并且访存模式可以用布局规划进行优化。适配度最高的包括 GEMM、矩阵分解类迭代算法、卷积类的滑窗运算、以及各种带稀疏结构的图计算。这些场景的共同点是，片上存储复用的收益极大，调度自由度也高。

不适合的场景也有几个：小矩阵单次调用，比如几 KB 的矩阵做一次乘法，调度和规划的开销可能超过计算本身；还有极端串行依赖的算法，每一步必须等前一步全部完成，这种 DAG 几乎没有并行宽度，KDA 发挥不了优势。遇到这种负载，老老实实用专门的稠密优化库可能更划算。

后续我打算做两件事。第一是把 KDA 的核大小和驻留策略交给自动调优器去搜索，目前固定 64×64 只是经验值，不同架构下最优值一定不同。第二是把 KDA 的依赖图和调度层扩展成多 GPU 版本，因为单卡内存始终是天花板，只有把多卡的显存和计算资源统一调度起来，才能继续往更大矩阵推进。

最后说一个我自己折腾这么久总结出来的体会：任何一个高性能计算方案，逻辑正确只是及格，真正的提升往往来自“数据待在合适的位置”这个朴素的道理。矩阵记忆听起来像是存储的事，实际上它决定的是一台机器能把算力用出来多少；KDA 做的无非就是把这件事用工程化的手段固定下来，让每次计算都不再做无用功。