Catlass（CANN Atlas）算子模板库是异构计算平台中实现高性能通用矩阵乘法（GEMM）的基石。它超越了传统 BLAS 库的范畴，通过 C++ 模板元编程技术，将硬件架构细节（如 Cube Unit 的位宽、L0 缓存大小、DMA 引擎控制）直接编码到算子定义中。Catlass 的核心目标是在 LLM 等计算密集型任务中，实现计算单元的绝对饱和和内存延迟的完全掩盖。

CANN 组织链接： https://atomgit.com/cann
Catlass 仓库链接： https://atomgit.com/cann/catlass

---

1. Tiling 策略的层次化构建与数据局部性深度优化

矩阵乘法性能的关键在于如何将全局内存访问转化为片上缓存的重复利用。Catlass 模板通过三层 Tiling 结构实现这一点。

1.1 三层 Tiling 结构：从全局到寄存器

Catlass 模板将 GEMM 映射到硬件的多个存储层级，每层对应不同的循环维度。

• Outer Loop (Global/Stream Level)：控制整体的 Batch 或 M/N/K 分块，这部分通常与 Runtime 的 Stream 调度相关联。
• Middle Loop (UB/L1 Level)：负责将数据划分到本地统一缓冲区（UB）。这些块的尺寸必须精确计算，以匹配 L1 缓存的容量，并且是 DMA 预取的基本单位。
• Inner Loop (Cube Unit Level)：这是最关键的一层。块尺寸直接与 Cube 单元的计算精度（如 $16\times 16$ 或 $32\times 16$ 的计算模式）和寄存器堆大小相关联，确保计算核心在一次调度中完成最大量的乘累加（MAC）操作。

1.2 访存流水线与 DMA 引擎的精细同步

Catlass 模板中嵌入了精确控制数据预取的逻辑。

• 双缓冲机制的显式实现：在内层循环中，数据加载与计算操作被显式地交错。当 Cube Unit 处理当前 Tile $T_i$ 时，DMA 引擎被调度加载下一个 Tile $T_{i+1}$ 的数据。模板中的同步点保证了计算不会早于数据就绪。
• 地址指针的循环迭代：模板代码管理着 L0/L1 缓存的读写指针，这些指针的更新逻辑与 Cube Unit 的计算顺序严格同步，保证了数据在片上缓存中的有效性和时效性。

2. 融合算子（Fusion Kernels）的片上数据流控制

Catlass 不仅优化 GEMM 本身，更重要的是将 GEMM 与后续的向量操作融合，实现数据的“零回写”。

2.1 消除中间结果的全局内存周转

融合操作是减少内存带宽压力的最有效手段。

• Bias Add 的内联：对于 $C=A\cdot B+Bias$，Bias 向量被视为一个常数或在 UB 中加载的小型张量。它直接在 Cube Unit 的累加器完成后的读出阶段与累加结果进行相加，结果直接用于下一阶段的激活函数。
• 精度边界处理：融合的关键在于处理精度边界。例如，如果 GEMM 使用 FP16 计算并 FP32 累加，融合的激活函数必须能接收 FP32 的累加结果，并按需将其量化或转换为下一阶段所需的格式。

2.2 激活函数调用的向量化调度

融合后的激活函数直接利用了 Vector Unit 的高吞吐能力。

• 累加结果的直接输入：GEMM 计算的输出（通常在寄存器堆中）直接作为 Vector Unit 的输入，执行如 $\text{ReLU}(x)=\max (0,x)$ 或更复杂的 $\text{GeLU}$ 近似计算。
• 原子操作：这些向量操作在 Catlass 模板内部被编译为原子指令，与 GEMM 的执行流紧密绑定，整个 $A\cdot B+Bias+\text{Activation}$ 过程无需任何全局内存写入。

3. 模板元编程与编译期优化机制

Catlass 模板库的强大源于其编译期（Compile-Time）的智能决策能力，而非运行时的动态调度。

3.1 维度参数化与模板特化

• 代码生成机制：开发者通过设置模板参数（如矩阵的 M/N/K 维度、数据类型 $T$、是否转置 $\text{Transpose}$），编译器会生成一个针对该特定配置高度优化的核函数实例。
• 消除运行时分支：所有关于数据布局、Tiling 尺寸和 DMA 预取参数的逻辑都在编译时确定，最终生成的代码块中不包含任何运行时分支判断，极大地提高了执行速度和确定性。

3.2 硬件特性与指令集的精确映射

Catlass 代码通过对特定硬件特性的直接引用，实现了指令的精确控制。

• Cube 模式选择：模板中存在对特定硬件模式的枚举。例如，选择支持 $16\times 16$ 整数 MAC 模式，或者 $8\times 8$ BF16 浮点模式，编译器据此生成相应的底层机器码。

Catlass 模板实例化过程的概念代码结构如下：

// 模板定义的核心骨架（概念性表示）
template <typename T_A, typename T_B, typename T_C, size_t M_BLOCK, size_t N_BLOCK, size_t K_BLOCK, bool TRANS_A>
__global__ void GemmTemplateKernel(...) {
    // 1. DMA 预取指令：加载 A_tile, B_tile 到 L0/UB
    DataLoad<T_A, TRANS_A>(A_global_ptr, ...);
    DataLoad<T_B>(B_global_ptr, ...);
    
    // 2. 核心计算循环（深度嵌套）
    for (size_t i = 0; i < M_BLOCK; ++i) {
        for (size_t j = 0; j < N_BLOCK; ++j) {
            // 触发 Cube Unit 硬件乘累加
            CubeAccumulate<T_A, T_B, T_C>(A_tile[i], B_tile[j], C_accumulator);
        }
    }
    
    // 3. 融合后处理（如 Bias Add）
    PostProcessFuse(C_accumulator, Bias_vector, Output_C); 
    
    // 4. 结果写回（或直接用于下一个融合算子）
    DataStore<T_C>(C_accumulator, C_global_ptr); 
}

4. 多精度管理与数值精度维护的精细化控制

Catlass 模板必须在实现 FP16/BF16 高速计算的同时，保证最终结果的数值可接受性。

4.1 累加器与精度提升的策略

• 累加精度提升：在所有低精度（FP16/INT8）的 GEMM 实现中，Catlass 模板确保中间累加器（Accumulator）采用更高的精度（通常是 FP32 格式）。这有效防止了在累加大量小数值时，因精度损失导致的数值下溢问题。

4.2 精度边界的显式处理

• 溢出/下溢检查：在将高精度累加结果转回目标低精度格式（如 FP16）时，模板代码包含逻辑来检测和处理数值溢出（> Max FP16）或下溢（接近零），确保输出值落在目标格式的有效范围内。

5. 性能验证与生态集成点

ops-math 提供了原子操作，而 Catlass 提供了完整的 GEMM 解决方案。GE 和 Runtime 负责调度这些 Catlass 实例。

5.1 编译产物与 Runtime 的交互

Catlass 编译出的核函数被打包进 OM 文件，Runtime 负责实例化。

• 参数固化：所有 Tiling 尺寸、数据类型、是否转置等信息都被 GE 固化到 OM 文件中，Runtime 直接加载这些参数，无需运行时重新计算 Tiling 策略。

5.2 性能调优的关键指标

• Cube Pipe 饱和度：在 Profiling 报告中，关注 Cube 单元的吞吐量指标。理想情况下，吞吐量应接近理论峰值，这表明内存搬运和计算调度已实现最大重叠。
• DMA 延迟掩盖率：评估数据预取操作（DMA）的完成时间与计算开始时间的重叠百分比。高重叠率意味着 Catlass 模板的异步控制是成功的。

---

CANN 组织链接： https://atomgit.com/cann
Catlass 仓库链接： https://atomgit.com/cann/catlass