如果你在昇腾 NPU 上做算子开发，迟早会碰到一个问题：写一个高性能的 GEMM（矩阵乘法）算子，到底有多难？

答案是：非常难。要手动做 tiling、手动管理 SRAM、手动处理 DMA 搬运、手动做指令流水……光是这些底层细节就能让一个算子开发者掉一层皮。

catlass 就是为了解决这个问题而存在的。它是昇腾官方出的算子模板库，专门给 Ascend C 开发者提供一套经过验证的高性能算子模板。名字里有个 “c” 开头，很多人第一反应是"这就是昇腾版的 CUTLASS 吧"——没错，但也不完全对。今天把这个问题说清楚。

---

catlass 到底是什么

catlass 是昇腾 CANN 开源社区里的一个仓库，定位是算子开发的模板库。它提供了一系列经过高度优化的算子实现模板，开发者不需要从零开始写 tiling 和 SRAM 管理逻辑，直接基于模板做适配就行。

这里有个关键区分要讲清楚：catlass 不是 CUTLASS 的封装，也不是 CUTLASS 的昇腾移植版。CUTLASS 是 NVIDIA 的 CUDA 模板库，针对的是 CUDA 编程模型和 GPU 硬件；catlass 针对的是 Ascend C 编程模型和昇腾达芬奇架构。两者解决的问题类似（都是降低高性能算子开发门槛），但底层完全不一样。

如果你看到有人说"catlass 就是 CUTLASS"，这个"就是"要打个大大的问号。正确理解是：catlass 在昇腾生态里的角色，类似于 CUTLASS 在 CUDA 生态里的角色——都是模板库，但实现和目标硬件完全不同。

昇腾异构计算架构里，catlass 的位置在这里：

Ascend C（算子编程语言）
  └─ catlass（算子模板库，提供可复用模板）
      └─ ops-nn / ops-math / ops-blas（基于 catlass 开发的算子库）
          └─ ATB / graph-autofusion（推理加速层，调用上面的算子）

---

catlass 解决了什么痛点

在 catlass 出现之前，昇腾 NPU 上的算子开发流程是这样的：

1. 用 Ascend C 写算子逻辑
2. 手动设计 tiling 策略（怎么分块、块多大、怎么装进 SRAM）
3. 手动管理 DMA 搬运（数据在 DDR、SRAM、L1 之间的搬移）
4. 手动做指令流水（把计算和数据搬运 overlap 起来）
5. 反复调优，看 profiler 找瓶颈

这个过程里，第 2、3、4 步是纯体力活——GEMM 的 tiling 策略是有套路的，SRAM 管理也是有最佳实践可以参考的，但每个开发者都要自己从头写一遍，效率极低。

catlass 把这三步变成了选模板、填参数：

// 没有 catlass：从零写 GEMM 算子，要手写 tiling + SRAM 管理
// 下面只是示意，实际 Ascend C 代码量大概是这个的 10 倍

// 手动算 tiling
int block_m = 128, block_n = 128, block_k = 32;
int num_blocks_m = (M + block_m - 1) / block_m;
// ... 几十行 tiling 逻辑

// 手动管理 SRAM
__aicore__ void Compute(int block_idx) {
    // 手动 DMA 搬运
    // 手动 SRAM 分配
    // 手动指令流水
}

// 有 catlass：基于模板，只需要配置参数
#include "catlass/catlass_gemm_template.h"

// 配置模板参数
using GemmConfig = catlass::GemmTemplate<
    /* DataType    = */ float16,
    /* BlockM      = */ 128,
    /* BlockN      = */ 128,
    /* BlockK      = */ 32,
    /* UseL2Cache  = */ true
>;

// 实例化模板，直接生成高性能 GEMM 算子
catlass::GemmKernel<GemmConfig> gemm_op;
gemm_op.Invoke(...);  // 调用，底层 tiling/SRAM/流水全部自动处理

核心差异：原来一个月的开发量，用 catlass 模板可能一周就搞定了，而且性能还不差——因为 catlass 的模板是昇腾官方团队调优过的，比你手写的第一个版本大概率要好。

---

catlass 的核心模板：覆盖哪些算子类型

catlass 目前提供的模板主要集中在矩阵计算和卷积类算子，具体包括：

GEMM 系列：标准矩阵乘法（C = A × B），支持 float16、bfloat16、float32 多种精度。这是最常用的模板，大模型里的 Linear 层底层全靠这个。

Batch GEMM：批量矩阵乘法，一次算多组小矩阵乘法。Transformer 里的 QKV 投影就可以用这个模板加速。

Conv 系列：卷积算子模板，支持 2D/3D 卷积，适配 CV 类模型（YOLO、ResNet 等）。

Element-wise 模板：逐元素操作（ReLU、GELU、LayerNorm 等）的高性能实现模板。这些操作看起来简单，但要做好 SRAM 利用率和指令流水，手写还是挺麻烦的。catlass 的模板帮你把这件事搞定了。

每一个模板都内置了针对昇腾达芬奇架构的优化策略：

• 自动 tiling：根据 SRAM 大小和矩阵尺寸，自动选择最优分块大小
• 双缓冲 DMA：计算和搬运 overlap，隐藏数据搬运延迟
• 指令调度优化：最大化 MTE（矩阵计算引擎）的利用率

---

catlass 和 CUTLASS 的详细对比

既然很多人把 catlass 和 CUTLASS 放在一起比，那我就认真比一下，看看异同在哪里。

维度	CUTLASS (NVIDIA)	catlass (昇腾)
目标硬件	NVIDIA GPU (CUDA)	昇腾 NPU (Ascend C)
编程模型	CUDA C++ 模板	Ascend C 模板
核心优化	Tensor Core 利用	达芬奇架构 MTE 利用
支持精度	fp16/bf16/fp32/int8	fp16/bf16/fp32（逐步补齐）
上手门槛	高（要懂 CUDA + Tensor Core）	中（要懂 Ascend C）
性能上限	非常高（NVIDIA 官方调优）	高（昇腾官方调优）

相同点：都是模板库，都解决了"手写高性能算子太难"的问题，都提供了 GEMM/Conv 等核心算子的优化模板。

核心不同点：CUTLASS 深度绑定 NVIDIA 的 Tensor Core 指令集，catlass 深度绑定昇腾达芬奇架构的 MTE（矩阵计算引擎）。你不能用 CUTLASS 的模板去写 Ascend C 算子，反过来也不行。

如果你是从 CUDA 生态迁移过来的开发者，catlass 的编程体验和 CUTLASS 有点像——都是"选模板、配参数、生成高性能算子"。但底层的硬件语义完全不同，不能指望"学会了 CUTLASS 就能直接上手 catlass"。Ascend C 的编程模型和 CUDA 还是有本质差异的，这个要心里有数。

---

catlass 怎么用：一个完整示例

假设你要开发一个自定义的矩阵乘法算子，用 catlass 模板来做，流程是这样的：

第一步：选模板。 去 catlass 仓库里找对应算子类型的模板。GEMM 就选 catlass_gemm_template.h，卷积就选 catlass_conv2d_template.h。

第二步：配参数。 根据你的数据类型、矩阵尺寸、性能需求，配置模板参数：

// 示例：配置一个 fp16 GEMM 模板
#include "catlass/catlass_gemm_template.h"

// 模板参数：数据类型、分块大小、是否用 L2 Cache
using MyGemm = catlass::GemmTemplate<
    catlass::DataType::FP16,   // A/B/C 数据类型
    128, 128, 32,              // BlockM, BlockN, BlockK
    true                        // 使用 L2 Cache 加速
>;

第三步：实例化调用。 把模板实例化成可调用的算子对象，在你的 Ascend C 代码里直接调用：

// 在 Ascend C kernel 里调用 catlass 模板
extern "C" __global__ __aicore__ void my_gemm_kernel(
    __gm__ half* a, __gm__ half* b, __gm__ half* c,
    int M, int N, int K
) {
    // 实例化 catlass GEMM 模板
    catlass::GemmKernel<MyGemm> gemm;
    
    // 调用，底层 tiling/DMA/流水全部自动处理
    gemm.Run(a, b, c, M, N, K);
}

第四步（可选）：调优。 如果默认模板参数的性能不够，可以根据你的具体矩阵尺寸，调整 BlockM/BlockN/BlockK 这几个分块参数。catlass 仓库里有一份调优指南，给出了不同矩阵尺寸下的最优参数配置参考。

---

几个踩坑经验

坑一：模板参数不是随便配的。 BlockM/BlockN/BlockK 这几个参数要匹配昇腾 NPU 的 SRAM 大小。配错了编译不会报错，但运行的时候性能会很差，甚至 SRAM 溢出。建议先用 catlass 提供的自动调优脚本跑一遍，找到适合你矩阵尺寸的参数组合，再写进代码。

坑二：catlass 模板不是万能的。 如果你的算子有特殊的融合需求（比如 GEMM + Softmax + Scale 三合一），catlass 的单一模板覆盖不到。这种情况可以基于 catlass 的 GEMM 模板做二次开发，把融合逻辑加进去——但这就不是"直接用模板"了，需要你懂 Ascend C 的融合算子开发。

坑三：catlass 依赖 opbase。 catlass 底层依赖 opbase 仓库里的基础组件（数据类型定义、SRAM 管理原语等）。如果你在编译的时候报了一堆"找不到 xxx 符号"的错误，先检查 opbase 有没有正确安装并加入 include 路径。

---

结尾

catlass 的核心价值是把算子开发从"手写汇编级优化"变成"选模板配参数"，大大降低昇腾 NPU 上高性能算子开发的门槛。

对于大多数算子开发者来说，不需要从零开始写 tiling 和 SRAM 管理——catlass 模板库里已经有经过官方调优的实现，直接用就行。如果你有特殊的融合需求或者性能目标，再基于模板做二次开发，也比从零开始效率高得多。

catlass 和 ops-nn、ops-math、ops-blas 这些算子库的关系是：后者是基于前者开发的成品算子，前者是后者的开发工具和模板基础。想深入理解昇腾算子生态，这两个层都要看。

源码在 https://atomgit.com/cann/catlass