• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend

本篇以 asc-devkit 仓库中的matmul_format_gemv样例为蓝本，讲解如何在 Ascend C 中实现 GEMV（General Matrix-Vector，矩阵向量乘）：当 Matmul 计算的 M=1 时，通过同时在 Tiling 侧与 Kernel 侧将矩阵 A 的数据格式设置为CubeFormat::VECTOR开启 GEMV 模式，使 1×K 的行向量按 16×16 分形一次性搬运进 L0A Buffer，显著减少搬运开销。读完本篇，你将掌握 GEMV 的开启方式、样例中 Kernel/Tiling/Host 三端关键实现，以及如何在本机完成编译、运行与精度校验的完整流程。

GEMV 模式下的矩阵乘计算原理，参考仓库 API 文档中的示意图：

一、样例概述：什么是 GEMV 模式

在标准 Matmul 计算中，左矩阵 A 的形状为 (M, K)，右矩阵 B 的形状为 (K, N)。当M=1时，该运算退化为 1×K 的行向量与 K×N 矩阵的乘法，即 GEMV 场景。

从 GEMV 特性文档 可以看到，开启 GEMV 模式的核心优势在于减少搬运数据量：

• GEMV 模式：将 A 矩阵从 L1 Buffer 搬运到 L0A Buffer 时，1×256 的向量被当作 16×16 的矩阵（即一个 16×16 输入分形）来处理，一次asc_copy_l12l0a调用即可完成 16×16 分形大小的矩阵搬运；
• 非 GEMV 模式：1×256 的向量被当作非对齐矩阵数据处理，需要将 M 方向对齐到 32 字节后搬运，要覆盖 K/16=16 个输入分形，搬运数据量相应增加。

两种模式的 M=1 矩阵乘计算对比可参考下图：

GEMV 模式还有两条需要了解的硬件/软件约束（引自上述特性文档）：

1. 1×K 矩阵 A 需满足512B 地址对齐，K 个数据连续存储。以 half 类型为例，当 K=256 时，软件侧可视作 16×16 分块，配置 m=1 后硬件解析为 1×256 向量，可通过asc_copy_l12l0a接口将 256 个 half 数据从 L1 Buffer 搬至 L0A Buffer；
2. C 矩阵是一个 (1×N) 的向量，当 1×N 向量被划分为多个 1×16 子向量时，每个子向量在 L0C Buffer 中占用 1024B，但实际有效数据仅占最低 64B。例如 N=50 时共划分 4 个 1×16 子向量，占用 4×1024B=4096B，而逻辑有效数据仅为 50×4B=200B。

本样例正是通过 Ascend C 高级 Matmul API 以声明式方式开启该能力：A 矩阵的 Format 参数设置为CubeFormat::VECTOR。在仓库头文件 matmul_config.h 中，CubeFormat枚举明确定义了各物理数据布局格式，其中VECTOR注释为 "1D array format for vector data"（一维数组格式，用于向量数据）：

enum class CubeFormat { ND = 0, // Standard format, row-major NZ, // Colum-major between blocks, row-major within blocks ZN, // row-major between blocks, Colum-major within blocks ZZ, // row-major between blocks, row-major within blocks NN, // Colum-major between blocks, Colum-major within blocks ND_ALIGN, // Aligned ND format, output aligned to 32-byte boundaries along N SCALAR, // A single numerical value with no dimensions VECTOR, // 1D array format for vector data ROW_MAJOR = ND, // Row-major format (equivalent to ND format) COLUMN_MAJOR = 8, // Column-major format (DN format) };

同样的枚举在 Tiling 侧也有一份定义，见 matmul_tiling_base.h。Kernel 侧与 Tiling 侧的 Format 必须保持一致，这是开启 GEMV 模式的两个必要条件（另一个是 M=1）。

二、支持的产品与 CANN 版本

本样例支持的产品及 CANN 软件版本如下：

对应的编译架构参数为：dav-2201对应 Atlas A2/A3 系列产品，dav-3510对应 Ascend 950PR&950DT 系列产品。

三、目录结构

样例位于 examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_format_gemv，目录组织如下：

├── matmul_format_gemv │ ├── scripts │ │ ├── gen_data.py // 输入数据和真值数据生成脚本文件 │ │ └── verify_result.py // 真值对比文件 │ ├── CMakeLists.txt // 编译工程文件 │ ├── data_utils.h // 数据读入写出函数 │ ├── matmul_format_gemv.asc // Ascend C样例实现 & 调用样例 │ └── README.md // 样例说明文档

四、样例规格

本样例的计算规格为M = 1, N = 32, K = 256，核函数名为matmul_format_gemv_custom：

注意一个细节：虽然逻辑上 A 是"行向量"，但其物理存储格式仍为 ND（行主序连续存储），满足 GEMV 要求的"K 个数据连续存储 + 512B 对齐"约束；CubeFormat::VECTOR描述的是 Cube 单元按向量方式解释和搬运该数据，而不是改变 GM 中的物理布局。

五、样例实现：三端关键代码剖析

样例文件 matmul_format_gemv.asc 是一个"Kernel + Tiling 生成 + Host 调用"三合一的单文件实现。以下按 Kernel、Tiling、Host 三个层面展开。

5.1 Kernel 关键步骤：Matmul 对象的 A 矩阵 Format 为 VECTOR

样例用模板类MatmulKernel封装了整个核函数逻辑。最关键的声明是Matmul对象本身——四个模板参数依次对应 A、B、C、Bias 的位置/格式/数据类型，其中 A 的 Format 即为CubeFormat::VECTOR：

AscendC::Matmul< AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::VECTOR, AType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, CType>, AscendC::MatmulType<AscendC::TPosition::GM, CubeFormat::ND, BiasType>, CFG_NORM> matmulObj;

核函数入口matmul_format_gemv_custom是 Cube 核函数（__cube__ __global__ __aicore__），流程为：

__cube__ __global__ __aicore__ void matmul_format_gemv_custom( __gm__ uint8_t* a, __gm__ uint8_t* b, __gm__ uint8_t* bias, __gm__ uint8_t* c, __gm__ uint8_t* workspace, __gm__ uint8_t* tilingGm) { if ASCEND_IS_AIV { return; // GEMV 由 Cube 核完成，向量核直接返回 } TCubeTiling tiling; CopyTiling(&tiling, tilingGm); AscendC::TPipe pipe; MatmulKernel<half, half, float, float> matmulKernel; REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, &tiling); matmulKernel.Init(a, b, bias, c, tiling); matmulKernel.Process(&pipe); }

从源码结构看，实现上有几个值得注意的设计点：

1. REGIST_MATMUL_OBJ注册调用：将matmulObj、系统 Workspace 与 Tiling 注册到 TPipe，使高级 Matmul API 能够接管数据搬运与 Cube 计算的流水调度，这是高级 API 相对于基础 API 的核心简化点；
2. CopyTiling辅助函数：Host 端生成的 Tiling 数据以字节流形式放在 GM 中，核函数内按uint32_t粒度循环拷贝回本地TCubeTiling结构，代码见 matmul_format_gemv.asc；
3. 多核偏移计算：Init中通过CalcOffset(AscendC::GetBlockIdx(), ...)依据 Tiling 结果把 M/N 切分到的多个核，每个核计算自己负责的 A/B/C/Bias 数据偏移；Process开头还会判断GetBlockIdx() >= tiling.usedCoreNum直接返回，处理尾块（tailM/tailN）时调用matmulObj.SetTail(tailM, tailN)。在本样例中 M=1，切分只发生在 N 方向；
4. 计算主体：Process中先SetTensorA/SetTensorB设置输入（附带IS_TRANS_A/IS_TRANS_B转置标志，本样例均为 false），可选SetBias，然后matmulObj.IterateAll(cGlobal)完成迭代计算，最后matmulObj.End()收尾。

5.2 Tiling 关键步骤：SetAType 同样设置为 VECTOR

Host 端的GenerateTiling函数使用matmul_tiling::MultiCoreMatmulTilingAPI 生成 Tiling 数据。与 Kernel 侧呼应，A 矩阵的 Format 同样传入matmul_tiling::CubeFormat::VECTOR：

void GenerateTiling(platform_ascendc::PlatformAscendC* ascendcPlatform, uint8_t* tilingBuf) { optiling::TCubeTiling tilingData; matmul_tiling::MultiCoreMatmulTiling tilingApi(*ascendcPlatform); tilingApi.SetDim(ascendcPlatform->GetCoreNumAic()); tilingApi.SetAType( matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::VECTOR, matmul_tiling::DataType::DT_FLOAT16, IS_TRANS_A); tilingApi.SetBType( matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT16, IS_TRANS_B); tilingApi.SetCType(matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); tilingApi.SetBiasType( matmul_tiling::TPosition::GM, matmul_tiling::CubeFormat::ND, matmul_tiling::DataType::DT_FLOAT); tilingApi.SetOrgShape(M, N, K); tilingApi.SetShape(M, N, K); tilingApi.EnableBias(IS_BIAS); tilingApi.SetBufferSpace(-1, -1, -1); int64_t res = tilingApi.GetTiling(tilingData); // Get matmul tiling data. if (res == -1) { std::cout << "gen tiling failed" << std::endl; return; } uint32_t tcubeTilingSize = tilingData.GetDataSize(); tilingData.SaveToBuffer(tilingBuf, tcubeTilingSize); }

SetAType的函数签名定义于 matmul_tiling_base.h：int32_t SetAType(TPosition pos, CubeFormat type, DataType dataType, bool isTrans = false)，即"位置 + 格式 + 数据类型 + 是否转置"四要素。SetDim传入核数为 -1 时的自动切分由 API 内部决定，最终tilingData通过SaveToBuffer序列化到 GM，供核函数使用。

5.3 Host 端调用实现：内核调用符 <<<>>>

main函数演示了完整的 ACL 调用链路：

1. 平台获取：通过platform_ascendc::PlatformAscendCManager::GetInstance()获取平台对象；CPU 调试模式下按架构返回模拟平台名（如Ascend910B1、Ascend950PR_9589）；
2. Tiling 生成：malloc(sizeof(TCubeTiling))后调用GenerateTiling；
3. ACL 初始化：aclInit→aclrtSetDevice→aclrtCreateContext→aclrtCreateStream；
4. 数据准备：aclrtMallocHost+ReadFile（读./input/x1_gm.bin、./input/x2_gm.bin）+aclrtMemcpy(ACL_MEMCPY_HOST_TO_DEVICE)将输入与 Tiling 数据搬上设备，输出与 Workspace 用aclrtMalloc在设备侧分配；
5. 内核启动：使用内核调用符<<<>>>以numBlocks = tiling->usedCoreNum个 Block 启动核函数：

matmul_format_gemv_custom<<<numBlocks, 0, stream>>>( aDevice, bDevice, biasDevice, cDevice, workspaceDevice, tilingDevice); aclrtSynchronizeStream(stream);

6. 结果回收：aclrtMemcpy(ACL_MEMCPY_DEVICE_TO_HOST)拷回 C 矩阵并WriteFile("./output/output.bin")，最后按相反顺序释放全部设备/主机内存并aclFinalize。

六、数据生成与精度校验脚本

6.1 gen_data.py：输入与真值数据生成

scripts/gen_data.py 与核函数规格保持完全一致（m, n, k, is_bias = 1, 32, 256, False）：

• 在(-1, 1)区间均匀随机生成x1_gm（shape [1, 256]）与x2_gm（shape [256, 32]），均转为float16存储到./input/x1_gm.bin、./input/x2_gm.bin；
• 将两者转为 float32 后做矩阵乘得到真值，写入./output/golden.bin（float32）；
• 若is_bias=True，则额外生成 bias 并参与真值计算（本样例IS_BIAS = false）。

6.2 verify_result.py：真值对比

scripts/verify_result.py 用np.isclose以rtol=atol=1e-3逐元素对比output/output.bin与output/golden.bin，统计误差元素占比error_ratio：

• 占比<= 1e-3（ERROR_TOL）时打印test pass!并返回 0；
• 否则打印前 100 个差异元素（index、期望值、实际值、相对误差）以及完整数据，进程以退出码 1 结束。

输入文件读取依赖 data_utils.h 中的ReadFile/WriteFile，前者基于stat+ifstream校验文件存在性与大小后二进制读入，后者以open(O_RDWR|O_CREAT|O_TRUNC)+write落盘。

七、编译与运行

以下操作在本样例根目录（examples/01_simd_cpp_api/04_advanced_api/00_matmul/matmul_format_gemv）下执行。

7.1 配置环境变量

根据当前环境上 CANN 开发套件包的安装方式，配置环境变量：

source ${install_path}/cann/set_env.sh

说明：${install_path}为 CANN 包安装目录，未指定安装目录时默认安装至/usr/local/Ascend下。

7.2 样例执行

在本样例目录下执行如下命令：

mkdir -p build && cd build; # 创建并进入build目录 cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式 python3 ../scripts/gen_data.py # 生成测试输入数据 ./demo # 执行编译生成的可执行程序，执行样例 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确

使用CPU 调试或NPU 仿真模式时，添加-DCMAKE_ASC_RUN_MODE=cpu或-DCMAKE_ASC_RUN_MODE=sim参数即可：

cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式 cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式

注意：切换编译模式前需清理 cmake 缓存，可在 build 目录下执行rm CMakeCache.txt后重新 cmake。

7.3 编译选项说明

这两个选项在 CMakeLists.txt 中声明为缓存变量，并通过target_compile_options将架构映射为编译器参数--npu-arch=${CMAKE_ASC_ARCHITECTURES}：

set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim") set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510") find_package(ASC REQUIRED) project(kernel_samples LANGUAGES ASC CXX) add_executable(demo matmul_format_gemv.asc) target_link_libraries(demo PRIVATE tiling_api # Matmul Tiling API（MultiCoreMatmulTiling） register # Matmul 对象注册（REGIST_MATMUL_OBJ 等） platform # 平台信息（PlatformAscendC） m dl )

从源码结构看，tiling_api、register、platform三个库恰好对应.asc文件中用到的三类能力：matmul_tiling::MultiCoreMatmulTilingTiling 生成、Matmul 对象的管道注册、PlatformAscendCManager平台对象获取。

7.4 执行结果

执行结果如下，说明精度对比成功：

test pass!

八、小结与扩展建议

• 开启 GEMV 的两个条件：M=1，且 A 矩阵在 Kernel 侧Matmul模板参数与 Tiling 侧SetAType中均声明为CubeFormat::VECTOR，两侧必须一致；
• 收益本质：把 1×K 向量按 16×16 输入分形整体搬运进 L0A，替代非对齐路径下逐分形、M 方向 32 字节对齐的多次搬运，降低搬运开销；
• 约束提醒：A 需 512B 对齐且 K 方向连续；L0C 侧每个 1×16 子向量占 1024B 而有效数据仅 64B，大 N 场景需留意 L0C 空间占用；
• 规格扩展：修改.asc中的constexpr uint32_t M/N/K与gen_data.py中对应常量（两者必须保持一致）即可复用到其他 GEMV 规格；
• 本样例属于 Matmul 高级 API 样例集 的一部分，同目录下还有matmul_format_column_major（列主序）、matmul_splitk（Split-K）、matmul_vecout（Vector 输出）等变体样例，可对比不同格式与切分策略下 Matmul 的配置差异。

• 人工智能
• 深度学习
• 算子库
• CANN
• Ascend