快速上手指南
环境准备
环境配套信息，可查看README中软件硬件配套说明。
下载CANN开发套件包，点击下载链接选择对应的开发套件包Ascend-cann-toolkit_<version>_linux-<arch>.run。 CANN开发套件包依赖固件驱动，如需安装请查阅安装NPU驱动固件页面。
安装CANN开发套件包。以下为root用户默认路径安装演示。
chmod +x Ascend-cann-toolkit_<version>_linux-<arch>.run
./Ascend-cann-toolkit_<version>_linux-<arch>.run --install
设置环境变量
source /usr/local/Ascend/ascend-toolkit/set_env.sh
使用CATLASS开发Matmul算子
本示例主要展示如何基于CATLASS快速搭建一个NPU上的BasicMatmul实现。示例中使用已提供的下层基础组件完成Device层和Kernel层组装，并调用算子输出结果。CATLASS分层示意图见api文档。
Kernel层算子定义
Kernel层模板由Block层组件构成。这里首先定义三个Block层组件。
<class BlockMmad_, class BlockEpilogue_, class BlockScheduler_>。
BlockMmad_为block层mmad计算接口，定义方式如下：
using DispatchPolicy = Catlass::Gemm::MmadAtlasA2Pingpong<true>; //流水排布使用
using L1TileShape = Catlass::GemmShape<128, 256, 256>; // L1基本块
using L0TileShape = Catlass::GemmShape<128, 256, 64>; // L0基本块
using AType = Catlass::Gemm::GemmType<ElementA, LayoutA>; //封装了A矩阵的数据类型和排布信息
using BType = Catlass::Gemm::GemmType<ElementB, LayoutB>; //封装了B矩阵的数据类型和排布信息
using CType = Catlass::Gemm::GemmType<ElementC, LayoutC>; //封装了C矩阵的数据类型和排布信息
using BlockMmad = Catlass::Gemm::Block::BlockMmad<DispatchPolicy,
L1TileShape,
L0TileShape,
AType,
BType,
CType>;
BlockEpilogue_为block层后处理，本文构建基础matmul，不涉及后处理，这里传入void。
using BlockEpilogue = void;
BlockScheduler_该模板类定义数据走位方式，提供计算offset的方法。此处使用定义好的GemmIdentityBlockSwizzle。参考Swizzle策略说明文档了解更多swizzle信息。
using BlockScheduler = typename Catlass::Gemm::Block::GemmIdentityBlockSwizzle<>;
- 基于上述组件即可完成BasicMatmul示例的Kernel层组装。
using MatmulKernel = Catlass::Gemm::Kernel::BasicMatmul<BlockMmad, BlockEpilogue, BlockScheduler>;
Device层算子定义
基于Kernel层组装的算子，完成核函数的编写。
- 使用CATLASS_GLOBAL修饰符定义Matmul函数，并传入算子的类型参数。
template <
class LayoutA,
class LayoutB,
class LayoutC
>
CATLASS_GLOBAL
void BasicMatmul(
GemmCoord problemShape,
GM_ADDR gmA, LayoutA layoutA,
GM_ADDR gmB, LayoutB layoutB,
GM_ADDR gmC, LayoutC layoutC);
- BasicMatmul的调用接口为
()运算符，需要传入Params作为参数。
typename MatmulKernel::Params params{problemShape, gmA, layoutA, gmB, layoutB, gmC, layoutC};
- 最后，实例化一个kernel，并执行该算子。
MatmulKernel matmul;
matmul(params);
算子调用
调用算子我们需要指定矩阵的输入输出的数据类型和数据排布信息，并使用<<<>>>的方式调用核函数。
BasicMatmul<<<BLOCK_NUM, nullptr, stream>>>(
options.problemShape, deviceA, layoutA, deviceB, layoutB, deviceC, layoutC);
算子编译
- 设置源代码对应的编译语言，可以混合使用纯C++代码和算子代码。对于算子文件，需要设定语言为ASCEND。
- 调用
catlass_example_add_executable函数指定target名称和编译文件。 - 如下所示，
00_basic_matmul为target名称，basic_matmul.cpp为需要编译的文件，dav-c220为NPU核心代号，一般填写dav-c220即可。
# CMakeLists.txt
set_source_files_properties(basic_matmul.cpp PROPERTIES LANGUAGE ASCEND)
catlass_example_add_executable(
00_basic_matmul
dav-c220
basic_matmul.cpp
)
在项目目录下，调用build.sh，即可编译examples中的kernel代码。
# 编译examples内所有用例
bash scripts/build.sh catlass_examples
# 编译指定用例
bash scripts/build.sh 00_basic_matmul
算子执行
切换到可执行文件的编译目录output/bin下，执行算子样例程序。
cd output/bin
# 可执行文件名 |矩阵m轴|n轴|k轴|Device ID（可选）
./00_basic_matmul 256 512 1024 0
执行结果如下，表明基于CATLASS编写的Kernel已经成功执行。
Compare success.
代码样例
完整的基础matmul样例参照examples/00_basic_matmul/basic_matmul.cpp。 该示例支持A/B矩阵为rowMajor数据排布输入。
版权声明
Copyright (c) 2025 Huawei Technologies Co., Ltd.
This file is a part of the CANN Open Software. Licensed under CANN Open Software License Agreement Version 1.0 (the "License"). Please refer to the License for details. You may not use this file except in compliance with the License.
THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. See LICENSE in the root of the software repository for the full text of the License.