昇腾CANN算子模板库catlass：从手写Ascend C到模板化开发的效率跃迁

前言

刚接触昇腾CANN算子开发的时候，我花了整整两天手写一个MatMul算子——Tiling逻辑、数据搬运、Cube单元调用，光是调通就折腾了一个下午。后来才知道，catlass这个仓库早就把这些重复劳动封装成了模板，三行配置就能生成一个经过优化的矩阵乘法算子。如果你也在做昇腾NPU上的算子开发，catlass能帮你省掉大量的底层搬运和编排工作，让你把精力放在业务逻辑而非硬件细节上。CANN生态里的catlass仓库，定位和NVIDIA CUTLASS类似，但专门为昇腾NPU的达芬奇架构做了适配，是昇腾算子编程语言Ascend C生态里不可或缺的模板库。

catlass是什么

catlass的全称是CANN Template Library for Ascend，它是昇腾CANN开源社区提供的算子模板库，对标NVIDIA的CUTLASS，但底层完全基于昇腾达芬奇架构的Cube和Vector计算单元来设计。

NVIDIA生态里的CUTLASS，做GPU编程的人基本都用过——它把GEMM（通用矩阵乘法）拆成流水线级的M、N、K分块，提供了一套可复用的模板框架，开发者只需要填入数据类型和分块参数就能得到一个高性能的矩阵乘法实现。catlass做的事情一样，只是目标硬件换成了昇腾NPU。

昇腾NPU和GPU在硬件设计上有本质区别。GPU靠大量线程并行，用SIMT模型；昇腾NPU用的是VLIW（超长指令字）+ 数据流架构，一个AI Core里面有Cube单元负责矩阵乘、Vector单元负责向量运算、Scalar单元负责标量控制。catlass的模板就是针对这种三单元协作的结构来设计的，不像CUTLASS那样基于线程块和共享内存来做流水线。

catlass目前提供三大类模板：

一是Matmul模板，覆盖FP16、FP32、BF16等数据类型的矩阵乘法，支持批量矩阵乘法和分组矩阵乘法。这是catlass最核心的部分，因为矩阵乘法占据了深度学习绝大部分的计算量。

二是Conv模板，提供卷积运算的模板化实现。卷积在昇腾上的实现思路是先把im2col展开，再走Matmul路径，catlass把这个流程封装成了一个完整的模板，开发者不需要自己处理im2col和分块的衔接。

三是Elementwise模板，覆盖逐元素运算，比如激活函数、Scale、BiasAdd等。这类算子计算密度不高，但调用频率很高，catlass通过Vector单元的模板化编排来减少数据搬运次数。

为什么要用catlass而不是手写Ascend C

手写Ascend C算子并不是做不到，而是投入产出比太低。以一个简单的FP16 MatMul为例，手写流程大概是这样的：

首先，你需要手动计算Tiling参数——把M、N、K三个维度切成AI Core本地内存能装下的小块，这个过程需要考虑L1 Cache和L0 Cache的容量限制，还要对齐到Cube单元的计算粒度（比如16x16的矩阵块）。

然后，你需要手动编排数据搬运流程——从Global Memory到L1，从L1到L0A/L0B，计算完再从L0C搬回Global Memory。每一次搬运都要写DataCopy指令，还要处理边界对齐和双缓冲（Double Buffer）来掩盖搬运延迟。

接着，你需要手动调用Cube单元的Matmul指令，设置M、N、K参数，处理矩阵的内存排布（行优先还是列优先）。

最后，你还需要考虑多核并行——把矩阵按行或按列分给不同的AI Core，处理核间同步和结果合并。

整个过程下来，一个简单的MatMul算子至少需要300行Ascend C代码，而且稍微改一下数据类型或者矩阵尺寸，Tiling参数就得重新算。

catlass把这个流程模板化了。你只需要提供矩阵的形状、数据类型和分块策略，catlass会自动生成Tiling计算、数据搬运编排和Cube调用的完整代码。来看一个具体的对比例子：

手写一个FP16 MatMul的Ascend C核函数，核心逻辑大致如下（省略了Tiling计算和边界处理）：

// 手写Ascend C的MatMul核心循环（简化版）// 这段代码只展示了单次分块计算，实际还需要外层循环遍历K维度extern"C"__global__ __aicore__voidmatmul_kernel(GM_ADDR x,GM_ADDR y,GM_ADDR z){// 定义Tile大小，这些参数需要手动计算和对齐// 为什么选16x16？因为Cube单元一次计算的最小粒度是16x16constexprintTILE_M=16;constexprintTILE_N=16;constexprintTILE_K=32;// 定义本地内存空间// L0A/L0B是Cube专用的输入buffer，L0C是输出buffer// 这里用双缓冲来掩盖数据搬运延迟__cbuf__ half tileA_buf[2][TILE_M*TILE_K];// 双缓冲L0A__cbuf__ half tileB_buf[2][TILE_K*TILE_N];// 双缓冲L0B__cbuf__floattileC_buf[TILE_M*TILE_N];// L0C，累加用FP32// 从Global Memory搬运第一个Tile到缓冲0// 这里省略了地址偏移计算，实际需要根据行列索引算出GM地址DataCopy(tileA_buf[0],x,TILE_M*TILE_K);DataCopy(tileB_buf[0],y,TILE_K*TILE_N);for(intk=0;k<K/TILE_K;k++){// 异步预取下一个Tile到另一个缓冲// 双缓冲的核心：计算当前块的同时搬运下一块intnext_buf=(k+1)%2;if(k+1<K/TILE_K){DataCopy(tileA_buf[next_buf],x+(k+1)*TILE_K,TILE_M*TILE_K);DataCopy(tileB_buf[next_buf],y+(k+1)*TILE_K*N,TILE_K*TILE_N);}// 调用Cube单元执行矩阵乘// MMAD是矩阵乘加指令，结果累加到L0Cintcur_buf=k%2;MMAD(tileC_buf,tileA_buf[cur_buf],tileB_buf[cur_buf],TILE_M,TILE_N,TILE_K);}// 把结果从L0C搬回Global Memory// FP32转FP16在这里做，避免单独再跑一次转换DataCopy(z,tileC_buf,TILE_M*TILE_N);}

这段代码只是最核心的循环部分，实际完整的实现还要加上Tiling参数计算、多核分发、边界处理、内存对齐检查，代码量至少翻三倍。而且这段代码只支持FP16，换成BF16就要改数据类型和Cube调用参数。

用catlass做同样的事情：

# catlass模板配置（Python接口）# 只需要声明矩阵形状、数据类型和分块策略importcatlass# 定义MatMul模板参数# 为什么用128x128的Tile？经验值，在Ascend 910上对L1 Cache利用率最高matmul=catlass.Matmul(M=1024,N=1024,K=1024,dtype_a=catlass.FP16,# 输入矩阵A的数据类型dtype_b=catlass.FP16,# 输入矩阵B的数据类型dtype_c=catlass.FP32,# 输出矩阵C用FP32累加，精度更高tile_m=128,# M方向分块大小tile_n=128,# N方向分块大小tile_k=64,# K方向分块大小double_buffer=True,# 开启双缓冲，掩盖搬运延迟)# 生成完整的Ascend C代码# catlass自动处理Tiling计算、数据搬运编排、Cube调用matmul.generate()# 输出到当前目录

catlass版本的配置代码不到20行，而且改数据类型只需要换一个参数，改矩阵尺寸也不需要重新计算Tiling。catlass内部会根据矩阵尺寸和AI Core数量自动计算最优的分块策略和核间分发方案。

catlass的内部架构

catlass的设计分三层，从上到下分别是用户接口层、模板引擎层和代码生成层。

用户接口层就是上面看到的Python API，负责接收用户的配置参数并做基本的合法性校验——比如检查Tile大小是不是16的倍数、数据类型组合是否支持。

模板引擎层是catlass的核心，它维护了一组预定义的计算流程模板。每个模板描述了一种计算模式（Matmul/Conv/Elementwise）在昇腾AI Core上的完整执行流程，包括数据搬运路径、计算指令序列、同步点和循环结构。模板引擎会根据用户的配置参数，实例化这些模板——把抽象的Tile大小替换成具体数值，把数据类型替换成具体的Ascend C类型声明。

代码生成层把实例化后的模板转换成可编译的Ascend C代码。这一层会处理一些底层细节，比如生成正确的内存地址计算代码、插入Pipeline Barrier指令确保数据一致性、生成多核并行的核函数入口。

整个流程的关键在于模板引擎层——它把硬件相关的优化决策（Tiling策略、双缓冲配置、Cube调用参数）封装在了模板内部，用户不需要了解达芬奇架构的细节就能得到一个高质量的实现。

catlass和CUTLASS的关键差异

虽然catlass对标CUTLASS，但两者在技术实现上有几个根本性的差异。

计算模型不同。CUTLASS基于GPU的SIMT模型，用线程块（Thread Block）和共享内存（Shared Memory）来实现分块计算和数据复用。catlass基于昇腾的VLIW模型，用AI Core的本地内存（L1/L0）和数据搬运指令来实现同样的目标。CUTLASS的流水线靠线程级并行来驱动，catlass靠指令级的双缓冲来驱动。

内存层次不同。GPU有Global Memory、Shared Memory、Register File三级，CUTLASS主要在Shared Memory这一级做分块和数据复用。昇腾NPU有Global Memory、L1 Cache、L0A/L0B/L0C三级，catlass在L1和L0这两级做分块，而且L0A/L0B是Cube单元专用的，不能像Shared Memory那样通用访问。

数据排布不同。CUTLASS默认使用行优先（Row-Major）排布，和cuBLAS的列优先排布需要做转换。catlass使用昇腾的5D排布格式（NC1HWC0），其中C0维度固定为16，对应Cube单元的计算粒度。这种排布格式在卷积场景下比行优先更高效，因为不需要做im2col展开后的额外重排。

流水线策略不同。CUTLASS的流水线是WMMMA指令驱动的，每个线程块内的多个Warp协作搬运和计算。catlass的流水线是DataCopy+MMAD指令驱动的，通过Double Buffer在搬运和计算之间做时间重叠，不需要多Warp协作。

使用前后的效率对比

用一个具体的例子来看。假设我们要实现一个FP16的GEMM，矩阵尺寸为4096x4096x4096，在Ascend 910上运行。

几个值得关注的点：

小矩阵场景下catlass的优势更明显。1024x1024的GEMM，手写版本很容易选到不合适的Tiling参数，导致Cube利用率偏低。catlass内部有一套基于矩阵尺寸的启发式规则，会自动选择更小的Tile大小来提高Cube利用率。

开发时间的差距是最大的。手写版本两到三天的开发时间，catlass只需要30分钟。而且catlass生成的代码质量是稳定的，不会因为开发者经验不足导致性能瓶颈。

性能上两者差距不大，因为手写版本在调优后也能达到接近硬件峰值的水平。但catlass的默认配置就能达到不错的性能，不需要反复调优。对于项目初期需要快速验证的场景，catlass的效率优势非常明显。

实际使用中的注意事项

catlass虽然好用，但有几个限制需要了解。

模板覆盖的算子类型有限。目前catlass主要覆盖Matmul、Conv、Elementwise三大类，如果你要实现的是Reduction、Sort这类不规则算子，catlass暂时没有对应的模板，还是得手写Ascend C。catlass的模板库在持续扩充，CANN 8.0版本新增了FlashAttention相关的模板，后续版本会继续扩展。

自定义Tiling策略的灵活性有限。catlass内部的Tiling启发式规则覆盖了常见的矩阵尺寸，但对于非常规尺寸（比如K维度远大于M和N的瘦长矩阵），默认Tiling可能不是最优的。这种情况下，你可以通过tile_m/tile_n/tile_k参数手动指定分块策略，但需要自己对硬件有足够的了解。

生成代码的可读性一般。catlass生成的Ascend C代码是模板实例化的结果，变量命名和代码结构和手写代码差别较大，如果需要二次修改，建议在catlass的模板层面修改而不是直接改生成代码。

性能对比的补充说明

上面的性能数据基于Ascend 910平台，使用CANN 8.0版本测试。不同硬件平台的性能会有差异——Ascend 950PR和950DT的AI Core数量、L1 Cache大小和Cube峰值算力都不同于Ascend 910，catlass会根据目标硬件自动调整Tiling参数。

对于卷积算子，catlass的Conv模板会比手写版本有更大的性能优势，因为卷积的im2col+Matmul流程比较复杂，手写时容易在im2col的内存排布上出问题，catlass把这部分逻辑封装在了模板内部。

结尾

catlass解决的核心问题是：让昇腾NPU上的算子开发从"手写底层指令"变成"配置模板参数"。它不是要取代Ascend C，而是在Ascend C之上提供了一层抽象，把重复的Tiling计算、数据搬运编排和Cube调用封装成了可复用的模板。对于需要频繁实现矩阵乘法、卷积等标准算子的开发者来说，catlass能显著缩短开发周期，同时保证生成代码的性能接近手写调优的水平。如果你的项目里有大量的标准算子需要实现，catlass值得认真看看。

https://atomgit.com/cann/catlass