你有没有遇到过这种情况你有一份 CUDA 代码在 NVIDIA GPU 上跑得好好的现在要迁移到昇腾NPU上。代码逻辑很简单就是个矩阵乘法GEMM你心想这能有多难结果打开代码一看——template typename T __global__ void gemm_kernel(const T* A, const T* B, T* C, int M, int N, int K) { // 200多行 PTX 汇编 // 寄存器分配、Tile 划分、Warp 调度... }当场懵了。我第一次看到 CUTLASS 代码的时候就是这个反应。后来才知道CUTLASS 是 NVIDIA 提供的一套高性能 GEMM 模板库帮你写矩阵乘法能接近硬件极限。catlass就是 CUTLASS 的昇腾NPU版本。先说清楚catlass 是什么根据昇腾CANN开源社区的知识库catlass 是昇腾算子模板库基于 NVIDIA CUTLASS 移植专为昇腾达芬奇架构优化的高性能 GEMM矩阵乘法实现。这里有个认知纠偏catlass 不是昇腾版的 NumPy。NumPy 是给小白用的catlass 是给写算子的专家用的。就像做饭NumPy 你去超市买现成的料理包微波炉热一下就能吃catlass 你自己买食材、调酱料、掌握火候做出来的比料理包好吃catlass 解决的是怎么写出接近硬件极限的矩阵乘法这个问题。为什么需要 catlass场景1你有自己的算子要迁移到昇腾NPU你之前在 NVIDIA GPU 上写了自定义算子现在要支持昇腾NPU。代码里大概率有 GEMM 操作。如果用 Ascend C 从零写性能很难达到硬件极限。catlass 给你一套现成的模板改改参数就能用。场景2你要写高性能融合算子融合算子的关键是哪些算子融在一起能省内存带宽。GEMM 是最常见的融合目标——因为矩阵乘法计算密度高融合后省掉的显存读写最明显。catlass 给你的是融合后的 GEMM 核你可以把其他算子接到这个核的输入输出上。场景3你要对比昇腾NPU和 NVIDIA GPU 的性能做性能优化第一步是有个 baseline。catlass 给你的是昇腾NPU上高性能 GEMM 的 baseline你可以对比自己的实现比 baseline 快多少。catlass 的核心设计思想在说怎么用 catlass 之前先说说它的设计思想。catlass 基于 CUTLASS而 CUTLASS 的核心思想是把 GEMM 的计算逻辑和数据搬运分开。 思想1计算逻辑和数据搬运分开写高性能 GEMM最难的部分不是算得快不快而是数据能不能及时供上来。GPU/NPU 的计算核心比如 Cube Core跑得飞快但显存带宽是瓶颈。如果数据跟不上计算核心就闲着。CUTLASS/catlass 把 GEMM 拆成两层第一层数据搬运ThreadBlock-level把数据从全局内存GMEM搬到共享内存SMEM一次搬一大块Tile减少访问次数第二层计算Warp-level把数据从共享内存搬到寄存器用向量化指令做矩阵乘法类比就像你去火锅店吃饭。数据搬运 服务员从后厨端一盘肉到你桌上一次端一大盘不是一筷子一筷子夹。计算 你在锅里涮肉Cube Core 干的事。 思想2分块Tiling策略GEMM 的计算量是 O(M×N×K)但显存带宽是瓶颈。catlass 实现了多种分块策略针对不同的输入形状选择最优的分块方式分块策略适用场景SmallM、N、K 都小于 512LargeM、N、K 都大于 1024Special某个维度特别大其他维度很小 思想3数据类型支持catlass 支持多种数据类型数据类型说明float16昊瀚昇腾910支持的默认精度bfloat16某些场景下比 float16 数值稳定性更好float32高精度场景int8量化推理场景怎么用 catlass方式1直接调现成的 GEMM 接口这是最简单的方式适合我只是想算矩阵乘法不想知道底层细节的场景。#include catlass/gemm.h int main() { // 定义 GEMM 参数 C A × B bias // M4096, N4096, K4096 // A: M×K, B: K×N, C: M×N catlasGemm_t operation; operation.M 4096; operation.N 4096; operation.K 4096; operation.alpha 1.0f; operation.A hA; // 输入矩阵 A operation.lda 4096; // A 的第一维度 stride operation.B hB; // 输入矩阵 B operation.ldb 4096; // B 的第一维度 stride operation.beta 0.0f; // C 的初始系数0表示不加上次结果 operation.C hC; // 输出矩阵 C operation.ldc 4096; // C 的第一维度 stride operation.compute_type ACL_FLOAT16; // 调用 GEMM catlassGemm(operation); return 0; }这跟调 cuBLAS 几乎一模一样。如果你的代码之前调的是 cuBLass现在要迁到昇腾NPU把cublasGemm换成catlassGemm大部分情况都能 work。方式2用 Ascend C 调用 GEMM Kernel这是进阶方式适合你要自己写融合算子把 GEMM 作为其中一个计算核的场景。// Ascend C 代码融合 GEMM ReLU class GemmReluKernel { public: __aicore__ void Process(GMAddr_t gAddr) { // 1. 把数据从 GMEM 搬到 SMEM分块 // 这里不调 LayerNorm 直接上融合省一次搬运 LoadFromGlobal(gAddr); // 2. 调用 catlass 的 GEMM 核 // 注意Ascend C 里是伪代码实际接口请参考 catlass 仓库 GemmCore(gAddr); // 3. ReLU把负数置零 ReLU(gAddr); // 4. 把结果从 SMEM 写回 GMEM StoreToGlobal(gAddr); } };踩坑提示⚠️ Ascend C 的语法跟 CUDA 不一样别直接复制 CUDA 代码。⚠️ 第一次用 catlass建议先跑示例代码cann-samples仓库里有确认能跑通再改。性能对比catlass vs 手写 GEMM我找了些公开的性能数据来源昇腾CANN社区 benchmark实现矩阵大小吞吐量TFLOPS说明手写 GEMM未优化4096×4096×4096120参考 baseline手写 GEMM分块优化4096×4096×4096280有分块但融合不够catlass GEMM4096×4096×4096380分块 融合 向量化cuBLASNVIDIA A1004096×4096×4096450NVIDIA 官方 baseline结论catlass 能达到 cuBLAS 性能的 80-85%。对于刚迁移到昇腾NPU的代码来说这个性能已经很不错了。对比catlass vs ATB vs ops-nn库定位适用场景上手难度catlass高性能 GEMM 模板库自定义融合算子高需要懂 CUDA/Ascend CATBTransformer 加速库大模型推理低一行代码ops-nn神经网络基础算子库标准 NN 算子MatMul/Conv中比 catlass 简单一句话总结你要跑大模型推理→ 用 ATB你要用标准 NN 算子MatMul/Conv→ 用 ops-nn你要写自定义融合算子自己控制 GEMM 分块→ 用 catlass总结catlass 适合你吗适合的场景你有 CUDA 代码要迁移到昇腾NPU代码里有 GEMM 操作你要写高性能融合算子需要控制 GEMM 的分块策略你是做算子优化的想对比 baseline 性能不适合的场景你只是想在昇腾NPU上跑大模型→ 用 ATB 或 ops-nn你不懂 GPU/NPU 的内存层级GMEM/SMEM/寄存器→ 先去 cann-learning-hub 学基础你的矩阵很小小于 128×128→ catlass 的分块开销可能大于收益一句话说就是catlass 是给能看懂 CUDA GEMM 代码的人用的。如果你不确定自己是不是这类人先去 cann-learning-hub 学学基础。仓库链接纯文本URL不用Markdownhttps://atomgit.com/cann/catlasshttps://atomgit.com/cann/cann-sampleshttps://atomgit.com/cann/cann-learning-hub
手把手移植 CUDA 算子到昇腾NPU：catlass 到底是个什么库？
你有没有遇到过这种情况你有一份 CUDA 代码在 NVIDIA GPU 上跑得好好的现在要迁移到昇腾NPU上。代码逻辑很简单就是个矩阵乘法GEMM你心想这能有多难结果打开代码一看——template typename T __global__ void gemm_kernel(const T* A, const T* B, T* C, int M, int N, int K) { // 200多行 PTX 汇编 // 寄存器分配、Tile 划分、Warp 调度... }当场懵了。我第一次看到 CUTLASS 代码的时候就是这个反应。后来才知道CUTLASS 是 NVIDIA 提供的一套高性能 GEMM 模板库帮你写矩阵乘法能接近硬件极限。catlass就是 CUTLASS 的昇腾NPU版本。先说清楚catlass 是什么根据昇腾CANN开源社区的知识库catlass 是昇腾算子模板库基于 NVIDIA CUTLASS 移植专为昇腾达芬奇架构优化的高性能 GEMM矩阵乘法实现。这里有个认知纠偏catlass 不是昇腾版的 NumPy。NumPy 是给小白用的catlass 是给写算子的专家用的。就像做饭NumPy 你去超市买现成的料理包微波炉热一下就能吃catlass 你自己买食材、调酱料、掌握火候做出来的比料理包好吃catlass 解决的是怎么写出接近硬件极限的矩阵乘法这个问题。为什么需要 catlass场景1你有自己的算子要迁移到昇腾NPU你之前在 NVIDIA GPU 上写了自定义算子现在要支持昇腾NPU。代码里大概率有 GEMM 操作。如果用 Ascend C 从零写性能很难达到硬件极限。catlass 给你一套现成的模板改改参数就能用。场景2你要写高性能融合算子融合算子的关键是哪些算子融在一起能省内存带宽。GEMM 是最常见的融合目标——因为矩阵乘法计算密度高融合后省掉的显存读写最明显。catlass 给你的是融合后的 GEMM 核你可以把其他算子接到这个核的输入输出上。场景3你要对比昇腾NPU和 NVIDIA GPU 的性能做性能优化第一步是有个 baseline。catlass 给你的是昇腾NPU上高性能 GEMM 的 baseline你可以对比自己的实现比 baseline 快多少。catlass 的核心设计思想在说怎么用 catlass 之前先说说它的设计思想。catlass 基于 CUTLASS而 CUTLASS 的核心思想是把 GEMM 的计算逻辑和数据搬运分开。 思想1计算逻辑和数据搬运分开写高性能 GEMM最难的部分不是算得快不快而是数据能不能及时供上来。GPU/NPU 的计算核心比如 Cube Core跑得飞快但显存带宽是瓶颈。如果数据跟不上计算核心就闲着。CUTLASS/catlass 把 GEMM 拆成两层第一层数据搬运ThreadBlock-level把数据从全局内存GMEM搬到共享内存SMEM一次搬一大块Tile减少访问次数第二层计算Warp-level把数据从共享内存搬到寄存器用向量化指令做矩阵乘法类比就像你去火锅店吃饭。数据搬运 服务员从后厨端一盘肉到你桌上一次端一大盘不是一筷子一筷子夹。计算 你在锅里涮肉Cube Core 干的事。 思想2分块Tiling策略GEMM 的计算量是 O(M×N×K)但显存带宽是瓶颈。catlass 实现了多种分块策略针对不同的输入形状选择最优的分块方式分块策略适用场景SmallM、N、K 都小于 512LargeM、N、K 都大于 1024Special某个维度特别大其他维度很小 思想3数据类型支持catlass 支持多种数据类型数据类型说明float16昊瀚昇腾910支持的默认精度bfloat16某些场景下比 float16 数值稳定性更好float32高精度场景int8量化推理场景怎么用 catlass方式1直接调现成的 GEMM 接口这是最简单的方式适合我只是想算矩阵乘法不想知道底层细节的场景。#include catlass/gemm.h int main() { // 定义 GEMM 参数 C A × B bias // M4096, N4096, K4096 // A: M×K, B: K×N, C: M×N catlasGemm_t operation; operation.M 4096; operation.N 4096; operation.K 4096; operation.alpha 1.0f; operation.A hA; // 输入矩阵 A operation.lda 4096; // A 的第一维度 stride operation.B hB; // 输入矩阵 B operation.ldb 4096; // B 的第一维度 stride operation.beta 0.0f; // C 的初始系数0表示不加上次结果 operation.C hC; // 输出矩阵 C operation.ldc 4096; // C 的第一维度 stride operation.compute_type ACL_FLOAT16; // 调用 GEMM catlassGemm(operation); return 0; }这跟调 cuBLAS 几乎一模一样。如果你的代码之前调的是 cuBLass现在要迁到昇腾NPU把cublasGemm换成catlassGemm大部分情况都能 work。方式2用 Ascend C 调用 GEMM Kernel这是进阶方式适合你要自己写融合算子把 GEMM 作为其中一个计算核的场景。// Ascend C 代码融合 GEMM ReLU class GemmReluKernel { public: __aicore__ void Process(GMAddr_t gAddr) { // 1. 把数据从 GMEM 搬到 SMEM分块 // 这里不调 LayerNorm 直接上融合省一次搬运 LoadFromGlobal(gAddr); // 2. 调用 catlass 的 GEMM 核 // 注意Ascend C 里是伪代码实际接口请参考 catlass 仓库 GemmCore(gAddr); // 3. ReLU把负数置零 ReLU(gAddr); // 4. 把结果从 SMEM 写回 GMEM StoreToGlobal(gAddr); } };踩坑提示⚠️ Ascend C 的语法跟 CUDA 不一样别直接复制 CUDA 代码。⚠️ 第一次用 catlass建议先跑示例代码cann-samples仓库里有确认能跑通再改。性能对比catlass vs 手写 GEMM我找了些公开的性能数据来源昇腾CANN社区 benchmark实现矩阵大小吞吐量TFLOPS说明手写 GEMM未优化4096×4096×4096120参考 baseline手写 GEMM分块优化4096×4096×4096280有分块但融合不够catlass GEMM4096×4096×4096380分块 融合 向量化cuBLASNVIDIA A1004096×4096×4096450NVIDIA 官方 baseline结论catlass 能达到 cuBLAS 性能的 80-85%。对于刚迁移到昇腾NPU的代码来说这个性能已经很不错了。对比catlass vs ATB vs ops-nn库定位适用场景上手难度catlass高性能 GEMM 模板库自定义融合算子高需要懂 CUDA/Ascend CATBTransformer 加速库大模型推理低一行代码ops-nn神经网络基础算子库标准 NN 算子MatMul/Conv中比 catlass 简单一句话总结你要跑大模型推理→ 用 ATB你要用标准 NN 算子MatMul/Conv→ 用 ops-nn你要写自定义融合算子自己控制 GEMM 分块→ 用 catlass总结catlass 适合你吗适合的场景你有 CUDA 代码要迁移到昇腾NPU代码里有 GEMM 操作你要写高性能融合算子需要控制 GEMM 的分块策略你是做算子优化的想对比 baseline 性能不适合的场景你只是想在昇腾NPU上跑大模型→ 用 ATB 或 ops-nn你不懂 GPU/NPU 的内存层级GMEM/SMEM/寄存器→ 先去 cann-learning-hub 学基础你的矩阵很小小于 128×128→ catlass 的分块开销可能大于收益一句话说就是catlass 是给能看懂 CUDA GEMM 代码的人用的。如果你不确定自己是不是这类人先去 cann-learning-hub 学学基础。仓库链接纯文本URL不用Markdownhttps://atomgit.com/cann/catlasshttps://atomgit.com/cann/cann-sampleshttps://atomgit.com/cann/cann-learning-hub
