CANN矩阵乘模板库catlass在LLM推理中的实战应用：昇腾NPU上GEMM算子白盒化组装与硬件特化性能优化深度指南

前言

大模型推理场景下，矩阵乘法几乎决定了整个系统的性能天花板。无论是Transformer中的自注意力计算，还是FFN层的全连接变换，GEMM（通用矩阵乘法）的身影无处不在。在昇腾NPU上跑推理的人，多半都遇到过这样的困境：自带的算子库性能不够极致，自定义算子开发门槛又高得离谱，两头不到岸。

catlass 正是为了解决这个矛盾而生的。这个仓库在昇腾CANN生态里扮演了一个独特角色：它不是直接提供黑盒算子，而是把GEMM的计算过程拆成可组装的模板单元，让开发者从"能不能用"切换到"怎么用更好"的维度。从本质上看，catlass 解决的是硬件能力与业务需求之间的适配问题——昇腾NPU的Cube计算单元吞吐很高，但如果数据排布和分块策略没对上，分分钟被内存带宽拖成龟速。

这篇文章围绕catlass在LLM推理中的实战应用展开。不会从"什么是矩阵乘法"讲起，默认读者对GEMM有基本认知。重点放在三个地方：catlass的模板化设计原理是什么、白盒化组装的策略该怎么用、以及在LLM推理场景下能拿到多少真实收益。中间穿插代码示例和性能对比，把纸上谈兵和工程落地之间的距离拉近一些。

第一章 catlass在CANN算子生态中的位置与核心定位

1.1 算子生态全景：为什么需要catlass

昇腾CANN的计算服务层提供了大量现成算子，ops-nn仓里有matmul，ops-blas仓里也有高性能GEMM实现，ops-transformer仓里还有FlashAttention这类融合算子。这套生态足够丰富，常规推理任务拿来直接用没什么问题。

但问题在于，LLM推理对矩阵乘法的需求是高度场景化的。不同模型结构里，矩阵的shape差异巨大——从attention中的批量短序列到大矩阵的FFN层，tile策略需要动态适配。batch size从1到64不等时，最优的分块大小可能差出好几倍。同时，推理阶段还有量化计算的强需求，INT8、FP8的矩阵乘法必须在保持精度的前提下榨出硬件的最大吞吐。这些需求用黑盒算子是很难精细控制的。

这里就出现了一个根本矛盾：预置算子追求的是通用性，LLM推理追求的是极致性能，两者之间的gap需要一个中间层来弥合。catlass就是这个中间层——它不是另一个黑盒，而是把GEMM的计算模板完全展开，让开发者可以针对具体场景做硬件特化。

从架构层级来看，catlass位于CANN的昇腾计算语言层，底层依赖opbase提供的基础数据结构，往上对接ops-blas和ops-nn等算子仓，最终服务于cann-recipes-infer这类推理框架适配方案。整个链条里，catlass是性能深水区，开发者在这里拿到的是完全透明的硬件控制能力。

1.2 catlass与CUTLASS的关系：一个易混淆的设计选择

提到catlass，很多熟悉NVIDIA生态的人会立刻联想到CUTLASS。确实，两者在名字和设计理念上有明显的继承关系。CUTLASS是NVIDIA提供的CUDA矩阵模板库，而catlass则是昇腾基于Ascend C重写的对应实现。

但重写不是简单的语言翻译。CUDA和Ascend C的编程模型有本质差异，线程组织方式、内存层次结构、指令发射模式都不相同。catlass在移植CUTLASS核心理念的同时，针对昇腾达芬奇架构做了大量适配。分块策略、向量加载宽度、寄存器分配方式，这些细节都跟Ascend NPU的硬件特性强绑定。拿CUTLASS的经验直接套用到catlass上，十有八九会踩坑。

更重要的区别在于生态定位。CUTLASS在NVIDIA生态里更多是作为学习材料和二次开发的起点，而catlass在CANN生态里承担了更重的实战角色——它直接连接了硬件能力与上层推理框架，是从算子开发走向推理落地的必经之路。这种定位差异决定了catlass的文档风格、示例完整度、API设计都更偏向工程实用性。

1.3 模板库的核心能力边界

catlass不是万能的。它的核心能力是把GEMM的计算流程拆解成若干可配置的阶段，让开发者按需组装。具体来说，catlass提供了以下几类模板组件：数据分块模板负责把大矩阵切分成适合在片上缓存的小块；计算核心模板定义了矩阵乘累加的底层实现；数据排布模板处理矩阵在内存中的存储格式转换；再就是各种结尾操作（epilogue）的模板，用于接激活函数、量化反量化等后处理。

这套设计带来的灵活性是有代价的。开发者在享受精细控制权的同时，也需要理解分块策略为什么这样设计、tile大小为什么选这个值、不同shape下参数该怎么调整。这些都是硬骨头，不是调几个参数就能搞定的事情。所以catlass的目标用户其实很明确：有一定算子开发经验、需要针对特定场景做深度优化的团队。纯业务开发者如果只是想在昇腾NPU上跑通模型，直接用ops-nn仓里的matmul更省事。

第二章 catlass模板化设计原理：GEMM的拆解与重组

2.1 GEMM为什么要拆

GEMM的标准定义是 C = alpha * A * B + beta * C，看起来简单。但当矩阵规模达到LLM推理的量级时，单次矩阵乘法涉及的计算量和内存访问量都极其庞大。以一个70B参数模型为例，单层FFN的矩阵乘法维度可能达到4096×11008×4096，矩阵A是4096×11008，B是11008×4096，输出C是4096×4096。这样的规模不可能一次性全部送入计算单元，必须分块处理。

问题在于，分块策略的选择直接影响计算效率。分块太小，片上数据复用次数少，内存带宽成为瓶颈；分块太大，片上缓存装不下，频繁换入换出同样拖累性能。更麻烦的是，昇腾NPU有不同的计算单元——Cube负责矩阵乘，Vector负责向量操作，两者的最佳分块参数往往不同步。找到一个在当前硬件上表现优秀的分块方案，本身就是一个需要大量实验的工程问题。

catlass的思路是把这个工程问题结构化。它把GEMM的计算过程抽象成若干阶段，每个阶段有明确的数据流和计算语义，为每个阶段提供多种实现方案。开发者的工作从"设计整个GEMM流程"变成了"为每个阶段选择合适的实现"，复杂度从O(N)降到了O(N/M)——虽然还是要做决策，但每个决策的上下文变小了。

2.2 两级分块结构：Block Tile与Thread Tile

catlass的模板化设计中，最核心的概念是两级分块：Block Tile和Thread Tile。这个分层来自CUTLASS的设计经验，但在昇腾NPU上的具体参数取值完全是针对Ascend硬件调优过的。

Block Tile对应的是在片上缓存（On-Chip Buffer）中保存的数据块大小。昇腾NPU的片上缓存空间有限，必须精心设计A、B两个输入矩阵的tile排布，确保计算过程中需要的数据尽可能长时间驻留在缓存里。Block Tile的宽高比需要跟矩阵的实际shape匹配——方阵和非方阵的最优配置完全不同。catlass为常见shape区间提供了参考配置，这些配置来自硬件测试团队的实测数据，比拍脑袋定参数可靠得多。

Thread Tile对应的是单个计算线程处理的数据量。昇腾NPU上的Ascend C编程模型里，线程是最小调度单位，每个线程有独立的寄存器资源和本地内存。Thread Tile太小，寄存器利用率低；太大，寄存器溢出到本地内存，性能又掉下来。这个平衡点同样跟硬件绑定很深，catlass模板里给出的是经过充分验证的经验值。

理解两级分块的意义在于，开发者可以根据自己的矩阵shape和精度要求，在这个框架内做细粒度调优，而不是从零摸索。相当于站在一个经过充分测试的起点上往前跑，而不是每次都从起跑线开始。

2.3 模板参数体系：数据类型的抽象层

catlass的模板设计大量使用了C++模板元编程，这让它能够用同一套代码支持多种数据类型和计算模式。最基础的模板参数是Element（元素数据类型），可以是FP32、FP16、BF16、INT8等。不同数据类型对应的计算指令不同，累加精度也不同，catlass的模板会根据Element类型自动选择对应的计算路径。

Accumulator类型是另一个关键参数。在低精度计算中，中间结果通常需要用更高精度的累加器来避免溢出，比如INT8输入用INT32累加。catlass模板里可以单独指定累加器的类型，让计算精度和内存占用之间达到最优平衡。

Epilogue是结尾操作的抽象接口。GEMM计算完成后，通常还需要做relu、sigmoid等激活函数，或者接一个偏置向量加法。这些后处理操作可以融合到GEMM核里一次性完成，省掉中间结果的存储和读取开销。catlass的epilogue模板把这些操作统一抽象出来，支持开发者自定义新的结尾操作，灵活度非常高。

2.4 调度策略与数据流组织

分块方案和模板参数只是静态配置，真正的计算过程需要调度策略来驱动。catlass里定义了几种不同的调度策略，对应不同的数据流组织方式。

最基础的是Naive调度，即完全按照分块顺序逐个处理。这种方式实现简单，但数据复用效率不高，因为A矩阵的某个块在被B矩阵的所有块使用完之后才被淘汰，在那之前如果B矩阵的块已经被用过一遍，数据其实可以继续复用。

更高效的调度策略会把A矩阵的块固定住，遍历完所有对应的B矩阵块后再切换到下一个A块。这种方式显著提升了片上数据的复用率，但也带来了编程复杂度的提升——需要管理多个数据块的生命周期，确保在需要的时候数据还在缓存里。

catlass为常见场景预置了经过优化的调度策略。开发者在大多数情况下不需要从零设计调度逻辑，直接选用跟场景最接近的预置策略，在预置策略的基础上做小范围调整。这种分层设计把复杂度封装在底层，暴露给上层的接口足够简洁。

第三章 白盒化组装策略：把GEMM算子捏成你想要的样子

3.1 白盒组装的本质：从参数调优到逻辑定制

大多数开发者接触catlass的起点是调整已有模板的参数——改个tile大小、换个数据类型、试一下不同的epilogue。这种调整方式的收益是立竿见影的，改个参数跑个benchmark，性能数字当场就能看到变化。

但参数调优有天花板。当调优空间被榨干之后，想进一步提升性能，就必须进入白盒组装的层面：直接修改模板内部的计算逻辑、数据加载路径、流水线排布。这种方式没有现成的参数可以改，需要理解模板代码的内部结构，针对性地做定制。

白盒组装的典型场景包括：针对特定shape的手工向量化、针对特殊数据排布格式的加载逻辑定制、以及把多个小算子融合进GEMM核心以减少片外内存访问。这些场景的参数调优空间几乎为零，必须改代码才能解决。

以attention计算中的Batched GEMM为例。标准Transformer里的attention需要做Q、K、V三个矩阵的乘法，Q乘K的转置产生attention score，score乘V产生输出。如果把这两个矩阵乘分开跑，中间结果必须写回显存再读出来，额外的数据搬运开销在长序列场景下非常可观。如果能把两个矩阵乘融合成一次kernel调用，中间结果全程驻留在片上，性能提升是数量级的。

这种融合不是改参数能搞定的，必须深入到模板内部，理解数据流如何组织、流水线如何排布、寄存器资源如何分配，重新设计计算路径。这就是白盒组装的真正含义——不是在黑盒外面调旋钮，而是把盒子打开，按自己的需求重新布线。

3.2 数据排布与加载路径的定制

昇腾NPU的内存层次结构对GEMM性能影响极大。Global Memory、On-Chip Buffer、Register三个层级之间的数据搬运是性能瓶颈的根源。数据排布格式决定了加载效率，向量化加载能充分利用内存带宽，非对齐或者跨行访问则会引发额外的访存开销。

catlass模板中默认的数据排布是RowMajor或ColMajor，这是最通用的格式。但实际业务中，数据来源多种多样——PyTorch的tensor默认是RowMajor，昇腾的算子可能用不同的排布格式，如果数据排布跟模板期望的不一致，每次加载前都需要做格式转换，这部分开销在高频调用场景下是致命的。

白盒组装的一个常见任务是定制数据加载路径。开发者可以根据数据来源的实际排布格式，在模板中插入专门的格式转换逻辑，把转换开销从每次计算前挪到数据准备阶段。这样在计算核心执行时，数据已经是最佳排布状态，访存效率直接拉满。

加载路径的另一个优化点是向量化宽度。昇腾NPU的向量指令支持不同宽度的数据打包，FP16可以一次加载8个，INT8可以一次加载16个。加载宽度越大，单位指令处理的数据越多，但同时对数据对齐的要求也越高。定制加载路径时，需要在向量化宽度和数据实际分布之间找到最优匹配点。

3.3 流水线排布与计算重叠

高性能GEMM实现里，流水线（Pipeline）是绕不开的话题。理想情况下，数据加载和计算应该完全重叠——一边用当前计算所需的数据，一边预取下一批数据，这样访存等待时间被计算填满，整体吞吐接近计算单元的理论上限。

但流水线设计有其复杂性。数据依赖关系决定了哪些操作可以并行，哪些必须串行。如果预取节奏没控制好，可能出现数据还没加载完计算就开始等，或者数据加载完但计算单元还在处理上一批，两种情况都是性能损失。

catlass的模板里实现了基础的流水线逻辑，但预置流水线的参数是针对"平均场景"调的。当矩阵shape偏离平均值较大时，默认流水线的效率会下降。这个时候需要进入白盒层面，重新设计流水线的深度和切换时机。

一个常见的策略是动态调整流水线深度。短矩阵用浅流水线，减少启动开销；长矩阵用深流水线，最大化计算重叠收益。实现这个策略需要在模板代码里加入shape判断逻辑，根据判断结果选择不同的流水线配置。这种程度的定制已经完全超出了参数调优的范畴，必须改代码。

3.4 融合扩展：从GEMM到自定义融合算子

catlass的epilogue机制为融合扩展提供了标准接口。标准GEMM的结尾操作相对固定——加偏置、激活函数、量化反量化。但实际业务需求五花八门，有些场景需要在GEMM结尾做自定义的归一化操作，有些场景需要接特殊的激活函数，还有些场景需要在结尾做多路输出的分流。

这些需求都可以通过扩展epilogue接口来实现。catlass定义了epilogue的抽象基类，开发者继承这个基类实现自己的结尾逻辑，把这个自定义epilogue塞进GEMM模板的对应位置。整个过程不需要改动GEMM核心的计算路径，只在结尾阶段插入自己的逻辑。

融合扩展的实际收益取决于融合带来的访存节省能否覆盖增加的逻辑开销。如果自定义epilogue的计算量很小但需要一次额外的全局内存访问，融合可能反而变慢。如果自定义epilogue的计算量适中且能省掉至少一次全局内存读写，融合的收益就非常可观。判断标准其实很简单：融合后kernel总耗时是否比分离调用更短。

第四章 LLM推理场景下的实战优化：代码与效果

4.1 场景选择：FFN层矩阵乘与Attention计算

LLM推理有两个典型的矩阵乘热点：FFN层的两个矩阵乘（Gate Projection和Up Projection合并后的FFN1，以及FFN2）和Attention计算中的QKV投影与Score计算。这两处的矩阵规模差异很大——FFN层的中间维度通常远大于hidden维度，而Attention计算的维度结构则相对扁平。

以LLaMA架构为例，FFN层的输入维度是hidden_size，输出维度是intermediate_size，通常intermediate_size是hidden_size的数倍。以7B模型为例，hidden_size=4096，intermediate_size=11008，这意味着FFN1的权重矩阵是4096×11008，FFN2是11008×4096。两个矩阵乘的shape完全不对称，而且中间结果（FFN1的输出）需要作为FFN2的输入。

这种不对称shape的矩阵乘在通用模板里很难拿到最优性能，因为通用模板通常针对方阵或接近方阵的shape做优化。11008×4098这样的窄长矩阵，最优分块策略和4096×4096的方阵完全不同。catlass的白盒组装在这里就有用武之地——可以在通用模板的基础上，针对FFN层的特定shape做定制化调优，包括调整Block Tile的宽高比、优化加载路径以适配窄长矩阵的数据访问模式、以及调整epilogue以适配SwiGLU激活函数的特殊结构。

4.2 基础调用：直接用预置模板

先看最直接的使用方式：选一个跟场景最接近的预置模板，改几个参数就跑起来。这种方式的门槛最低，适合快速验证可行性。

// 使用预置GEMM模板计算 FFN 层的第一个矩阵乘// A: (M, K) = (batch, 4096), B: (K, N) = (4096, 11008), C: (M, N) = (batch, 11008)// Element=half_t 对应 FP16，Accumulator=float_t 对应 FP32 累加器usingGemmKernel=typenamecatlass::arch::Gemm<catlass::arch::Shape<128,256,128>,// BlockTile: 128x256 的分块大小half_t,// 输入输出数据类型float_t,// 累加器数据类型catlass::arch::RowMajor,// C 和 A 矩阵排布格式catlass::arch::ColMajor// B 矩阵排布格式>::GemmIdentity;

这段代码选了一个128×256×128的分块配置，FP16输入输出，FP32累加器。参数选择不是拍脑袋定的，而是跟LLM推理中常见的矩阵shape做了匹配——128是M方向的常见batch对齐值，256是中间维度N方向的一个经验值，128是K方向的向量加载宽度基准。直接用这个配置跑起来，性能通常比ops-nn的通用matmul好一些，但距离极限还有差距。

WHY解释：选择128×256×128这个分块大小而不是更大的值，是因为11008这个维度是4096的倍数，256的分块能让片上缓存高效复用A矩阵的列块，避免频繁从Global Memory读取同一列的不同片段。如果分块太小，缓存命中率低；如果分块太大，A矩阵的列块可能在被完全使用前就被淘汰。

4.3 进阶组装：针对FFN层shape定制分块策略

在预置模板的基础上，根据FFN层的实际shape做定制化分块。核心思路是让Block Tile的N方向尺寸（256）与intermediate_size（11008）建立更紧密的倍数关系，减少边界处理的开销。

// 针对 LLaMA-7B FFN1 定制 GEMM 核// inter_size=11008, hidden_size=4096, 做 FP16 计算template<typenameGemmConf>structFFN1GemmLauncher{staticvoidlaunch(half_t*ptr_c,// 输出矩阵指针，shape: (batch, 11008)consthalf_t*ptr_a,// 输入矩阵指针，shape: (batch, 4096)consthalf_t*ptr_b,// 权重矩阵指针，shape: (4096, 11008)constintM,// batch * seq_len，通常是动态的constintK=4096,constintN=11008){// 根据 M 动态选择 thread block 数量，保持 SM 利用率constintthread_blocks=(M+GemmConf::kThreadBlockM-1)/GemmConf::kThreadBlockM;// K 方向按 64 分块迭代，每次处理 64 列的 A 矩阵数据constintk_blocks=(K+64-1)/64;// 调用底层核函数，这里省略了核函数的launch配置细节ffn1_kernel<<<thread_blocks,GemmConf::kThreadsPerBlock>>>(ptr_c,ptr_a,ptr_b,M,K,N);}};

这段代码演示了如何把预置GEMM核封装成适合FFN1场景的launcher。核心改动是引入了k_blocks的概念——K方向按64列分块迭代，每次只加载A矩阵的64列数据，这个宽度正好是Ascend Vector单元的一个高效处理粒度。通过让k_blocks跟K维度解耦，launcher可以适应不同的hidden_size，只需要改一次常量。

WHY解释：K方向固定按64分块而不是让分块大小可变，是为了保证每次迭代的计算量相对均衡。如果K不是64的倍数，第一块会处理余数部分，逻辑需要特殊处理但影响不大，统一的64粒度让片上缓存的管理更简单，因为每次进入计算核的数据量是固定的。

4.4 性能对比：通用算子 vs catlass定制核

在相同硬件环境下，用ops-nn仓的通用matmul算子和catlass定制GEMM核分别跑FFN层的矩阵乘，在不同batch size下的性能差异是显著的。

使用前 vs 使用后 效率对比

这个表格里的"基线水平""提升明显"是概括性描述，不对应具体数字。实际性能数字跟具体硬件型号、矩阵shape、CANN版本都有关系。文章里不做具体数字的捏造，读者在真实环境中测到的结果才是可信的。

从定性角度看，定制核的收益来源主要有三块：一是减少了边界处理和特殊情况判断的分支开销，二是通过更优的分块策略提升了片上数据复用率，三是通过融合epilogue省掉了中间结果的存储读写。这三块收益叠加在一起，在长序列场景下效果尤其明显。

4.5 量化场景下的特殊处理

LLM推理普遍采用INT8或FP8量化来降低显存占用和加速计算。量化后的矩阵乘法在catlass里有专门的模板支持，核心是在计算核里插入反量化逻辑，把INT8的输入还原成FP16或FP32再参与矩阵乘累加。

量化GEMM的关键是反量化的时机选择。方案一是先对A和B两个矩阵做完整反量化，生成临时FP16矩阵，送入FP16 GEMM核计算。这种方式实现简单，需要额外的临时显存存储反量化结果，如果显存紧张就不划算。

方案二是把反量化融合进GEMM核内部，在每次加载矩阵元素时即时反量化，累加到INT32累加器中。这种方式不需要临时显存，但每个元素的加载开销增加了（多了个反量化指令），整体收益取决于计算密度是否足够高来摊薄这个开销。

// 量化场景下的GEMM核配置示例usingQuantGemmConf=typenamecatlass::arch::Gemm<catlass::arch::Shape<128,128,128>,// 分块保持不变int8_t,// 输入是 INT8int32_t,// 累加器用 INT32catlass::arch::RowMajor,catlass::arch::ColMajor,catlass::epilogue::QuantizedEpilogue// 量化结尾，支持 per-tensor 或 per-channel 量化>;

这段配置选了一个对称的分块（128×128×128），因为量化场景下的数据排布更紧凑，对称分块能更好地适配量化后矩阵的内存布局。量化结尾支持per-tensor和per-channel两种粒度的反量化系数，前者实现更简单但精度略低，后者需要多传一个系数数组但精度更好。

WHY解释：量化结尾选择per-tensor还是per-channel，不是简单的精度越高越好。per-channel量化需要额外的系数数组访问，如果这个数组不在片上缓存里，每次访问都要去Global Memory读，开销可能抵消精度提升的收益。在实际工程中，需要实测两种方式的端到端性能再做选择。

第五章 模板参数调优的工程实践：避坑与经验

5.1 shape适配：从理论最优到工程可行

理论最优的分块策略可以通过数学建模推导出来——目标是让片上缓存的命中率最高、寄存器利用率最高、流水线停顿最少。但理论最优往往有一个前提：数据规模无限大、内存带宽无限宽、缓存策略无限智能。

工程场景里这些前提都不成立。矩阵的K维度可能是质数，batch size可能是变化的，显存空间可能被其他tensor占走了一部分。这些"不完美"条件会让理论最优的分块策略在实际执行时性能退化。

catlass的参数调优建议从"差不多能用"的分块开始，逐步微调。初始分块的选择可以参考仓库里提供的参考配置，这些配置是硬件测试团队在标准条件下跑出来的基准值，偏离基准越远，性能损失越大。在基准值附近微调时，建议用二分法做搜索——每次把参数改大改小，记录性能变化方向，沿着性能上升的方向继续搜索，直到性能曲线趋于平缓。

一个常见的坑是只测单个shape的性能就下结论。LLM推理里的矩阵shape是高度分布的，从极小的batch（比如prefill阶段的单条请求）到较大的batch（比如decode阶段的批处理），最优分块可能完全不同。调参时必须覆盖足够多样的shape区间，否则优化完一个场景反而劣化了另一个场景。

5.2 精度与性能的权衡

用catlass做深度优化时，精度问题总是会冒出来。FP16比FP32快但精度差，INT8比FP16更快但精度损失更明显。这种权衡没有标准答案，必须结合具体业务场景来判断。

判断标准其实很朴素：下游任务对数值误差的容忍度是多少。在图像生成场景里，少量数值误差可能只表现为像素级的噪点，人眼几乎察觉不到。但在某些精确计算场景里（比如某些科学计算任务），哪怕0.1%的误差都是不可接受的。大模型推理的中间计算通常比较宽容，FP16足够覆盖大多数场景。

catlass的模板设计支持混合精度的实验——可以在计算路径的不同阶段使用不同的精度，比如矩阵乘核心用FP16，结尾的激活函数用FP32。这种混合策略在某些场景下能达到精度和性能的最佳平衡点。实验方法是先跑通全FP16版本作为基准，逐步把非关键路径切换到低精度，同时监控端到端的精度指标（BLEU、Perplexity、accuracy等），找到一个精度损失可接受的上限点，在那个位置截断。

5.3 调试与诊断：数据正确性优先

写catlass自定义GEMM核最常见的问题不是性能差，而是计算结果错。这个问题在调优阶段尤其容易出现——改了一个参数发现性能涨了，但精度悄悄崩了，或者改了流水线配置发现特定shape下结果不对。

数据正确性验证必须在性能优化之前完成。建议的做法是准备一组标准测试用例，覆盖方阵、扁矩阵、长矩阵、小矩阵等不同shape，以及FP32、FP16、INT8等不同精度配置。每次修改代码后，用这套测试用例做完整的回归验证，只有全部用例通过才认为修改是安全的。

测试用例的设计也有讲究。随机数矩阵是基础测试，但不够——随机数的数值分布比较均匀，某些边界条件可能测不到。推荐用特定pattern的矩阵，比如单位矩阵、全1矩阵、阶梯矩阵，这些矩阵在特定shape下的计算结果是可预测的，可以做精确对比。如果catlass核的输出和PyTorch/Numpy在相同输入下的参考输出完全一致（允许浮点误差范围内的偏差），基本可以认为核函数实现是正确的。

第六章 catlass与其他CANN算子仓库的协同

6.1 catlass与ops-blas的关系

ops-blas仓提供的是开箱即用的高性能GEMM算子，这些算子内部很可能就是基于catlass开发的。对大多数用户来说，ops-blas已经足够好用——不需要理解分块策略、不需要知道寄存器分配，调用接口跟标准BLAS几乎一样。

catlass的角色是ops-blas的上游。当ops-blas的预置算子不能满足需求时（比如shape不在优化范围内，或者需要特殊的量化格式），开发者可以基于catlass开发自定义算子来覆盖这些case。这种分层设计让大多数用户不用操心底层细节，同时给有定制需求的开发者提供了足够的控制能力。

实际项目中，推荐的策略是用ops-blas做快速验证和产品化，用catlass做深度优化和产品化后持续迭代。先在ops-blas上跑通业务逻辑，确认方向是对的，切到catlass定制核做性能攻关。两套方案共用一套测试用例和benchmark框架，确保优化过程不会破坏正确性。

6.2 catlass与ATB的互补

ATB（Ascend Transformer Boost）是昇腾提供的大模型Transformer加速库，主要覆盖FlashAttention、LayerNorm、RoPE等Transformer核心算子的融合实现。catlass提供的是底层的GEMM计算能力，两者定位不同但可以协同使用。

在Transformer架构里，attention计算需要做QK^T的矩阵乘，这个矩阵乘的shape是 (seq_len, head_dim) × (head_dim, seq_len)，输出是 (seq_len, seq_len) 的attention score。seq_len在推理阶段通常是变化的——prefill阶段是固定的长序列，decode阶段是逐步延展的短序列。ATB的FlashAttention已经对这种shape变化做了优化，但极端长度的seq_len（比如32768甚至更长）可能超出ATB的优化范围。

在这种情况下，可以用catlass定制一个专门处理超长序列attention score计算的GEMM核，嵌入ATB的框架里使用。ATB负责整体的计算图编排和算子融合，catlass负责在ATB覆盖不到的角落里提供定制化的计算能力。这种协同方式既利用了ATB的成熟框架，又补足了边缘case的性能。

6.3 与上层推理框架的集成路径

LLM推理框架（比如vLLM、Text Generation Inference、或者厂商自研的推理引擎）通常通过算子抽象层来接入后端算子。接入catlass定制核的方式有两种：一种是把定制核封装成符合框架算子抽象接口的模块，框架通过统一的算子调用接口来触发；另一种是把定制核的kernel注册到CANN的算子注册表里，让上层框架通过算子名来调用。

第一种方式更灵活，但需要开发者自己维护接口适配层。第二种方式更标准，但需要定制核的参数和接口符合CANN的算子规范（包括输入输出shape描述、dtype声明等）。两种方式的取舍取决于定制核的使用范围——如果是项目内部专用，第一种方式更省事；如果是通用能力需要对外输出，第二种方式更规范。

无论哪种方式，都需要在推理框架的benchmark里验证catlass核的实际收益。纸上算出来的性能提升不一定是真实收益——框架层面的调度开销、内存分配开销、跨语言调用开销都可能吃掉底层kernel的优化收益。只有端到端实测才能验证真实收益的大小。

结尾

catlass的本质是一个高性能矩阵乘算子工厂，而不是一个性能开箱即用的黑盒。开发者在这里拿到的是对GEMM计算全流程的透明控制权——数据分块怎么切、寄存器怎么分配、流水线怎么排布，这些在黑盒算子里完全不可见的细节，在catlass里都可以按需定制。这种透明度的代价是学习成本和开发成本的双重提升，从投入到产出的链路比直接调用现成算子要长得多。

仓库链接：https://atomgit.com/cann/catlass