Ascend C Matmul GEMV 实战：用 CubeFormat::VECTOR 在 asc-devkit 中实现矩阵向量乘-北京尧图网络

人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载本篇以 asc-devkit 仓库中的matmul_format_gemv样例为蓝本讲解如何在 Ascend C 中实现 GEMVGeneral Matrix-Vector矩阵向量乘当 Matmul 计算的 M1 时通过同时在 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)。当M1时该运算退化为 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/1616 个输入分形搬运数据量相应增加。两种模式的 M1 矩阵乘计算对比可参考下图GEMV 模式还有两条需要了解的硬件/软件约束引自上述特性文档1×K 矩阵 A 需满足512B 地址对齐K 个数据连续存储。以 half 类型为例当 K256 时软件侧可视作 16×16 分块配置 m1 后硬件解析为 1×256 向量可通过asc_copy_l12l0a接口将 256 个 half 数据从 L1 Buffer 搬至 L0A BufferC 矩阵是一个 (1×N) 的向量当 1×N 向量被划分为多个 1×16 子向量时每个子向量在 L0C Buffer 中占用 1024B但实际有效数据仅占最低 64B。例如 N50 时共划分 4 个 1×16 子向量占用 4×1024B4096B而逻辑有效数据仅为 50×4B200B。本样例正是通过 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 模式的两个必要条件另一个是 M1。二、支持的产品与 CANN 版本本样例支持的产品及 CANN 软件版本如下产品CANN 软件版本Ascend 950PR950DT 系列产品 CANN 9.1.0Atlas A3 系列产品 CANN 9.0.0Atlas A2 系列产品 CANN 9.0.0对应的编译架构参数为dav-2201对应 Atlas A2/A3 系列产品dav-3510对应 Ascend 950PR950DT 系列产品。三、目录结构样例位于 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项目名称shape数据类型数据格式样例类型 (OpType)Matmul样例输入a[M, K]halfND样例输入b[K, N]halfND样例输出c[M, N]floatND核函数名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::VECTORAscendC::Matmul AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::VECTOR, AType, AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, BType, AscendC::MatmulTypeAscendC::TPosition::GM, CubeFormat::ND, CType, AscendC::MatmulTypeAscendC::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; MatmulKernelhalf, half, float, float matmulKernel; REGIST_MATMUL_OBJ(pipe, GetSysWorkSpacePtr(), matmulKernel.matmulObj, tiling); matmulKernel.Init(a, b, bias, c, tiling); matmulKernel.Process(pipe); }从源码结构看实现上有几个值得注意的设计点REGIST_MATMUL_OBJ注册调用将matmulObj、系统 Workspace 与 Tiling 注册到 TPipe使高级 Matmul API 能够接管数据搬运与 Cube 计算的流水调度这是高级 API 相对于基础 API 的核心简化点CopyTiling辅助函数Host 端生成的 Tiling 数据以字节流形式放在 GM 中核函数内按uint32_t粒度循环拷贝回本地TCubeTiling结构代码见 matmul_format_gemv.asc多核偏移计算Init中通过CalcOffset(AscendC::GetBlockIdx(), ...)依据 Tiling 结果把 M/N 切分到的多个核每个核计算自己负责的 A/B/C/Bias 数据偏移Process开头还会判断GetBlockIdx() tiling.usedCoreNum直接返回处理尾块tailM/tailN时调用matmulObj.SetTail(tailM, tailN)。在本样例中 M1切分只发生在 N 方向计算主体Process中先SetTensorA/SetTensorB设置输入附带IS_TRANS_A/IS_TRANS_B转置标志本样例均为 false可选SetBias然后matmulObj.IterateAll(cGlobal)完成迭代计算最后matmulObj.End()收尾。5.2 Tiling 关键步骤SetAType 同样设置为 VECTORHost 端的GenerateTiling函数使用matmul_tiling::MultiCoreMatmulTilingAPI 生成 Tiling 数据。与 Kernel 侧呼应A 矩阵的 Format 同样传入matmul_tiling::CubeFormat::VECTORvoid 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.hint32_t SetAType(TPosition pos, CubeFormat type, DataType dataType, bool isTrans false)即位置格式数据类型是否转置四要素。SetDim传入核数为 -1 时的自动切分由 API 内部决定最终tilingData通过SaveToBuffer序列化到 GM供核函数使用。5.3 Host 端调用实现内核调用符 main函数演示了完整的 ACL 调用链路平台获取通过platform_ascendc::PlatformAscendCManager::GetInstance()获取平台对象CPU 调试模式下按架构返回模拟平台名如Ascend910B1、Ascend950PR_9589Tiling 生成malloc(sizeof(TCubeTiling))后调用GenerateTilingACL 初始化aclInit→aclrtSetDevice→aclrtCreateContext→aclrtCreateStream数据准备aclrtMallocHostReadFile读./input/x1_gm.bin、./input/x2_gm.binaclrtMemcpy(ACL_MEMCPY_HOST_TO_DEVICE)将输入与 Tiling 数据搬上设备输出与 Workspace 用aclrtMalloc在设备侧分配内核启动使用内核调用符以numBlocks tiling-usedCoreNum个 Block 启动核函数matmul_format_gemv_customnumBlocks, 0, stream( aDevice, bDevice, biasDevice, cDevice, workspaceDevice, tilingDevice); aclrtSynchronizeStream(stream);结果回收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_gmshape [1, 256]与x2_gmshape [256, 32]均转为float16存储到./input/x1_gm.bin、./input/x2_gm.bin将两者转为 float32 后做矩阵乘得到真值写入./output/golden.binfloat32若is_biasTrue则额外生成 bias 并参与真值计算本样例IS_BIAS false。6.2 verify_result.py真值对比scripts/verify_result.py 用np.isclose以rtolatol1e-3逐元素对比output/output.bin与output/golden.bin统计误差元素占比error_ratio占比 1e-3ERROR_TOL时打印test pass!并返回 0否则打印前 100 个差异元素index、期望值、实际值、相对误差以及完整数据进程以退出码 1 结束。输入文件读取依赖 data_utils.h 中的ReadFile/WriteFile前者基于statifstream校验文件存在性与大小后二进制读入后者以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_ARCHITECTURESdav-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_MODEcpu或-DCMAKE_ASC_RUN_MODEsim参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # cpu调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU仿真模式注意切换编译模式前需清理 cmake 缓存可在 build 目录下执行rm CMakeCache.txt后重新 cmake。7.3 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201 对应 Atlas A2 系列产品和 Atlas A3 系列产品dav-3510 对应 Ascend 950PR950DT 系列产品这两个选项在 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 APIMultiCoreMatmulTiling register # Matmul 对象注册REGIST_MATMUL_OBJ 等 platform # 平台信息PlatformAscendC m dl )从源码结构看tiling_api、register、platform三个库恰好对应.asc文件中用到的三类能力matmul_tiling::MultiCoreMatmulTilingTiling 生成、Matmul 对象的管道注册、PlatformAscendCManager平台对象获取。7.4 执行结果执行结果如下说明精度对比成功test pass!八、小结与扩展建议开启 GEMV 的两个条件M1且 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_splitkSplit-K、matmul_vecoutVector 输出等变体样例可对比不同格式与切分策略下 Matmul 的配置差异。赞分享人工智能深度学习算子库CANNAscend【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址https://gitcode.com/cann/asc-devkit点击查看免费下载相关推荐Ascend C 高阶 API 矩阵乘实战用 Matmul API 在 asc-devkit 中实现 C A * B BiasAscend C 高阶 API 矩阵乘实战用 Matmul API 在 asc devkit 中实现 C A B Bias 本文以 asc devki人工智能深度学习算子库CANNAscendasc-devkit Ascend C 批量矩阵乘直调实践用 IterateBatch 高效处理 BSNGD 多 Batch Matmulasc devkit Ascend C 批量矩阵乘直调实践用 IterateBatch 高效处理 BSNGD 多 Batch Matmul 本文围绕 asc人工智能深度学习算子库CANNAscendAscend C 算子开发实战MxMatmul 量化矩阵乘法直调样例深度解析asc-devkitAscend C 算子开发实战MxMatmul 量化矩阵乘法直调样例深度解析asc devkit 导读本文以 asc devkit 仓库中 exampl人工智能深度学习算子库CANNAscend上一篇Flutter IntroViews-Flutter打造令人惊艳的Material Design应用引导页下一篇Material Shell源码注释规范提升代码可读性的实践创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考

Ascend C Matmul GEMV 实战：用 CubeFormat::VECTOR 在 asc-devkit 中实现矩阵向量乘

相关推荐

REFramework 完整指南：RE Engine 游戏模组框架与 Lua 脚本平台快速跑通

DeepFilterNet ONNX 导出与跨平台部署：从导出命令到三条集成路径的完整教程

CloudCompare点云处理实战：从激光雷达数据到OBJ建模全流程

更多资讯

让AI放心改代码：claude-delegator的read-only与workspace-write沙箱双模式安全机制详解

Scrapy入门实战：从创建项目到跑通第一个爬虫的全流程

Python进阶与并发编程07：CPU工作为什么要换进程，以及为什么换完可能更慢

MySQL COUNT函数全解析：三种形态、统计语义与性能优化

SpringBoot整合MyBatis与Thymeleaf实现CRUD实战

2026.9.29：使用 Milvus Operator 在 Kubernetes 中运行 Milvus

今日更新

NetSpot 实战：用信道分析与频宽优化解决 Wi-Fi 卡顿

C++ split() 字符串分割的五种实现与性能实测

R语言坐标轴范围设置：从xlim/ylim到coord_cartesian的完整指南

本周精选

交通拥堵预测毕设：基于Python+LSTM的车辆流量预测系统全流程解析

MATLAB/Simulink搭建10机39节点电力系统暂态稳定仿真实战

无人机航拍三维重建全流程：从SfM到网格生成的避坑指南

本月热门

我发现了一个新思路：用 Remotion + Claude Code 像写代码一样自动化生成短视频

Windows下 Codex 中 Chrome 和 Computer Use 插件不可用问题排查及解决参考方式：TaoToken 统一 Key 配置与验证

2026 大模型集体涨价：用 Python 做企业 Token 成本测算与选型避坑（附配置）

读完文章，想聊聊您的网站？