Ascend C SIMD C API 入门实战cann-samples 核函数直调、向量计算、矩阵乘、融合与 RegBase 编程样例解析【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samplescann-samples 仓库的 01_simd_cpp_api 目录集中提供了基于 Ascend C SIMD C API 的入门样例覆盖核函数直调、向量计算、矩阵乘计算、融合计算以及 RegBase 向量编程五大基础场景。本文以该目录的官方概览文档为主线结合仓库内的样例源码、构建脚本与运行说明逐层拆解每个样例的工程结构、编程范式与编译运行流程帮助开发者在 Ascend 硬件上快速上手 SIMD C API并掌握加载-计算-存储三段式流水、多核并行切分、队列式内存管理等核心编程模式。一、样例目录定位从零开始的 SIMD C API 学习路径01_simd_cpp_api 是 cann-samples 中 SIMD单指令多数据C API 的入门级样例集合。与仓库内更高阶的性能故事如 2_Performance 下的各类性能调优实战不同本目录的定位是帮助开发者理解样例工程结构、编译运行流程和常见编程模式聚焦跑通第一个算子这一核心目标。其覆盖的五类基础场景正好构成一条由浅入深的学习主线核函数直调用内核调用符在 NPU 上启动核函数理解最小可运行工程向量计算以 Add 算子为例对比两种不同的内存与同步管理机制矩阵乘计算接触 Cube 单元的 Matmul 计算对比基础 API 与高级 API 两种实现融合计算将 Matmul 矩阵乘与 LeakyRelu 激活函数融合减少中间结果搬运RegBase 向量编程寄存器基RegBase编程范式下的向量计算与 VFVector Function融合。样例总览根据官方概览文档本目录包含的样例及其功能描述如下目录名称功能描述00_quickstartHelloWorld 核函数实现01_add基于不同编程模式的向量计算实现02_matrix矩阵乘计算03_fusion_operation融合计算04_reg_compute基于 RegBase 编程的向量计算从源码结构看各子目录均包含.asc内核源码文件、CMakeLists.txt构建文件与README.md说明文档向量计算类样例还附带scripts/gen_data.py与scripts/verify_result.py数据生成与校验脚本每个样例都遵循统一的实现 构建 运行 验证工程范式可直接作为新算子开发的脚手架。二、00_quickstartHelloWorld 核函数直调00_quickstart 是整套样例的起点通过 hello_world 样例演示了核函数在 NPU 侧运行验证的基础流程使用内核调用符启动核函数核函数内通过printf打印输出结果。2.1 最小内核工程结构Samples/0_Introduction/01_simd_cpp_api/00_quickstart/hello_world/ ├── CMakeLists.txt // 编译工程文件 ├── hello_world.asc // Ascend C 样例实现 调用样例 └── README.md // 样例说明文档其中 hello_world.asc 完整呈现了 Ascend C 内核开发的最小闭环#include utils/debug/asc_printf.h #include acl/acl.h __global__ __vector__ void hello_world() { printf(Hello World!!!\n); } int main(int argc, char const* argv[]) { aclrtSetDevice(0); // 运行管理资源申请。 aclrtStream stream nullptr; aclrtCreateStream(stream); hello_world8, 0, stream(); aclrtSynchronizeStream(stream); aclrtDestroyStream(stream); aclrtResetDevice(0); // 销毁运行资源。 return 0; }这段代码体现了 SIMD C API 内核直调的三个关键要素__global__ __vector__修饰符声明这是一个在 AIVAI Vector核上执行的向量内核函数printf输出能力来自 asc_printf.h 头文件8, 0, stream内核调用符第一个参数8指定使用 8 个核并行执行第二个参数0表示不使用局部内存参数第三个参数指定 ACL Stream源码位置ACL 运行时管理aclrtSetDevice/aclrtCreateStream申请设备与流资源aclrtSynchronizeStream等待内核执行完成aclrtDestroyStream/aclrtResetDevice释放资源构成完整的资源生命周期。2.2 构建系统与编译选项CMakeLists.txt 展示了 SIMD C API 工程的构建方式通过find_package(ASC REQUIRED)引入 Ascend 编译工具链以project(kernel_samples LANGUAGES ASC)声明工程语言再用add_executable(demo hello_world.asc)将.asc内核源文件编译为可执行程序并通过--npu-arch${CMAKE_ASC_ARCHITECTURES}编译选项指定 NPU 架构。本目录所有样例统一支持以下两个编译选项选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2/A3 训练与推理系列产品dav-3510对应 Ascend 950PR/Ascend 950DT2.3 编译、运行与结果验证在样例根目录执行以下步骤默认 npu 模式mkdir -p build cd build; cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程默认 npu 模式 ./demo若使用 CPU 调试或 NPU 仿真模式追加对应参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU 仿真模式注意切换编译模式前需清理 build 目录下的CMakeCache.txt缓存文件再重新执行 cmake。执行成功后输出 8 行打印每行对应一个 AIVAI VectorCore 的打印结果[AIV Block 0/8] Hello World!!! [AIV Block 1/8] Hello World!!! [AIV Block 2/8] Hello World!!! [AIV Block 3/8] Hello World!!! [AIV Block 4/8] Hello World!!! [AIV Block 5/8] Hello World!!! [AIV Block 6/8] Hello World!!! [AIV Block 7/8] Hello World!!![AIV Block X/8]表示当前为第 X 个核共 8 个核该输出证明核函数已在 8 个 AIV Core 上成功并行执行——这是理解后续多核并行计算样例每个核通过block_idx切分数据的基础。三、01_add向量计算的两种编程范式01_add 以 Add 算子z_i x_i y_i逐元素相加为载体验证 SIMD C API 的向量计算能力核心价值在于对比两种不同的内存与同步管理机制目录编程范式说明add静态 Tensor 编程使用LocalMemAllocator直接管理 UB 内存手动插入PipeBarrier同步add_tpipe_tqueTQue / TPipe 队列式编程通过EnQue/DeQue队列接口管理内存与流水同步两个样例的输入输出规格一致x、y为float类型、shape[8, 2048]、ND 布局输出z与输入同规格均以 8 核并行完成计算每核处理 2048 个元素blockLength 2048总计 8×204816384 个 float 元素。3.1 静态 Tensor 范式三段式流水线add 样例的计算逻辑遵循加载Load-计算Compute-存储Store三段式流水结构加载将输入数据 x、y 从 GMGlobal Memory片外全局内存容量大但访问慢搬运到 UBUnified BufferAI Core 片内向量计算专用缓存容量小但访问快经GlobalTensor与LocalTensor访问计算在 UB 上对xLocal与yLocal执行向量加法结果存入zLocal存储将计算结果从 UB 写回 GM。其核心实现涉及四个关键 API/变量DataCopyGM 与 UB 之间的数据搬运接口搬运方向由参数顺序决定PipeBarrierPIPE_ALL()流水线同步屏障确保数据搬运完成后才能执行后续操作避免读写冲突block_idx内置变量表示当前核的索引等价于GetBlockIdx()用于多核并行时的数据切分LocalMemAllocatorHardware::UBUB 片上内存分配器为向量计算分配连续的LocalTensor缓冲区。核心代码示意如下template uint32_t blockLength __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { AscendC::InitSocState(); // Global Tensor在 GM 上按 block_idx 偏移切分数据段实现多核并行 AscendC::GlobalTensorfloat xGm, yGm, zGm; xGm.SetGlobalBuffer(x block_idx * blockLength, blockLength); yGm.SetGlobalBuffer(y block_idx * blockLength, blockLength); zGm.SetGlobalBuffer(z block_idx * blockLength, blockLength); // Local Tensor在 UB 上分配计算缓冲区 AscendC::LocalMemAllocatorAscendC::Hardware::UB ubAllocator; AscendC::LocalTensorfloat xLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat yLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat zLocal ubAllocator.Allocfloat, blockLength(); // 阶段 1 加载GM - UB AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保加载完成后才能计算 // 阶段 2 计算z x y AscendC::Add(zLocal, xLocal, yLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保计算完成后才能存储 // 阶段 3 存储UB - GM AscendC::DataCopy(zGm, zLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保存储完成 }调用方式为numBlocks, 0, stream其中numBlocks8指定 8 核并行执行。整段实现以InitSocState()初始化 AI Core 硬件状态为起点SetGlobalBuffer(x block_idx * blockLength, ...)让每个核处理不同数据段三次PipeBarrierPIPE_ALL()分别保障加载→计算、计算→存储、存储完成的数据一致性。3.2 TQue/TPipe 范式队列式内存管理add_tpipe_tque 展示另一种编程思路——用 TPipe 和 TQue 统一管理内存分配与流水同步这是更接近生产级算子的写法。其处理流程为add_custom作为内核入口接收totalLength通过GetBlockNum()计算当前块需处理的数据长度用GetBlockIdx()计算当前核在 GM 中的起始位置用DataCopy将输入从 GM 搬入 UB并通过EnQue将输入LocalTensor放入输入队列通过DeQue从输入队列取出张量在 UB 上执行Add再通过EnQue将结果LocalTensor放入输出队列通过DeQue从输出队列取出结果用DataCopy写回当前核负责的 GM 数据段。该样例工程还包含完整的数据验证链路scripts/gen_data.py生成输入与真值数据scripts/verify_result.py校验输出与真值的一致性运行流程为mkdir -p build cd 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 # 校验输出结果与算法逻辑校验通过时输出test pass!。3.3 从 Add 到性能优化可优化方向分析官方文档还从实战角度给出了 Add 算子的四个可优化方向对理解后续性能样例极具启发序号优化方向当前实现问题预期收益1多核动态分配固定使用 8 核未按实际可用核数动态分配动态获取可用核数充分利用多核并行降低端到端时延2增大搬运粒度每次搬运 2048 个 float8KB粒度偏小增大单次搬运数据量、减少搬运次数摊薄启动开销提升带宽利用率3双缓冲流水并行加载、计算、存储严格串行MTE2/V/MTE3 硬件单元无法同时工作采用 Ping-Pong 双缓冲机制使三段并行、隐藏搬运时延4L2 Cache 旁路输入数据仅读取一次却默认经过 L2 Cache增加 Cache 污染对流式访问数据设置 L2 Cache 旁路减少无效 Cache 开销提升搬运效率四、02_matrix矩阵乘计算的三种 API 实现02_matrix 将场景从向量计算扩展到矩阵乘计算演示了 SIMD C API 面向 Cube 单元的三种实现路径目录实现方式支持产品matmul_basic_api基于静态 Tensor 编程范式实现矩阵乘950PR/950DT、Atlas A2/A3 训练与推理系列matmul_advanced_api基于 Matmul 高级 API 实现矩阵乘950PR/950DT、Atlas A2/A3 训练与推理系列matmul_tensor_api基于静态 Tensor API 编程范式实现矩阵乘950PR/950DT三个样例matmul_basic_api 与 matmul_advanced_api 均含scripts/gen_data.py、scripts/verify_result.py与 data_utils.h 数据工具头文件共享相同的数据生成与精度校验流程差异体现在内核侧 API 层次基础 API要求开发者显式管理 GlobalTensor/LocalTensor 与数据搬运Matmul 高级 API将切块、搬运与 Cube 计算封装为高层接口显著降低开发门槛静态 Tensor API则提供另一套静态张量编程范式。开发者可据此在同一计算目标下对比不同 API 的表达力与灵活性。五、03_fusion_operationCube 与 Vector 融合计算03_fusion_operation 引入自定义融合计算场景将 Matmul 矩阵乘与 LeakyRelu 激活函数融合为单算子避免中间结果落回 GM 带来的额外搬运开销目录功能描述支持产品matmul_leakyrelu_advanced_api基于高级 API 实现 Matmul 矩阵乘与 LeakyRelu 激活函数的融合计算950PR/950DT、Atlas A2/A3 训练与推理系列matmul_leakyrelu_basic_api基于基础 API 实现 Matmul 矩阵乘与 LeakyRelu 激活函数的融合计算950PR/950DT、Atlas A2/A3 训练与推理系列从实现角度看融合算子内部呈现典型的Cube矩阵乘 Vector激活函数协同流水Cube 单元完成矩阵乘后在片上直接衔接 Vector 单元的 LeakyRelu 计算两个样例分别验证了在基础 API 与高级 API 两种编程范式下完成跨单元融合的可行性是理解算子融合收益减少片内外数据往返、降低端到端时延的入门范例。六、04_reg_computeRegBase 向量编程04_reg_compute 介绍 SIMD C API 中RegBase寄存器基编程的向量计算能力目前仅支持 Ascend 950PR/Ascend 950DT目录功能描述add基于 RegBase 编程接口的 Add 示例gelu基于 LocalMemAllocator 与 RegBase/VF 融合的 GELU 计算示例RegBase 编程将数据与计算直接锚定在寄存器/寄存器组层面相比默认的内存基MemBase编程可获得更精细的数据布局与计算控制。其中 gelu 样例进一步将 RegBase 与 VFVector Function融合使用配合LocalMemAllocator完成片上内存分配展示了 RegBase 在非线性激活函数场景下的组合编程能力。对于希望深入寄存器级向量优化的开发者这是通往仓库内 Reg数据搬运场景选型指南 与 RegBase 性能故事如 softmax_regbase_story、kv_rms_norm_rope_cache_story的桥梁。七、贯穿全目录的调试与性能分析方法所有 SIMD C API 入门样例共享同一套调试与性能分析手段官方在 add 文档中给出了完整说明7.1 功能调试printf 与 DumpTensorAscendC::printf在内核侧需要输出日志的位置调用例如AscendC::printf(add blockIdx%d\n, AscendC::GetBlockIdx());适用于 CPU/NPU 域调试AscendC::DumpTensor用于 Dump 指定LocalTensor的内容并支持打印附加的自定义信息仅支持 uint32_t 类型例如在Add计算后调用AscendC::DumpTensor(zLocal, 1, 32);打印前 32 个元素。注意printf/DumpTensor 打印功能对算子实际运行性能有一定影响通常仅在调试阶段使用可通过设置ASCENDC_DUMP0关闭打印功能。7.2 性能调试msOpProf 单算子性能分析工具msOpProf 是单算子性能分析工具支持msopprof与msopprof simulator两种模式可对不同运行模式设备端/仿真和文件类型可执行程序/算子二进制.o文件进行性能数据采集与自动解析。设备端采集直接度量算子在 Ascend AI 处理器上的真实执行时间运行msopprof ./demo即可。命令执行后默认目录下生成OPPROF_{timestamp}_XXX文件夹其关键性能数据文件包括├──dump # 原始性能数据用户无需查看 ├──ArithmeticUtilization.csv # Cube/Vector 指令周期占比 ├──L2Cache.csv # L2 Cache 命中率影响 MTE2需合理规划数据搬运逻辑 ├──Memory.csv # UB、L1、主存的读写带宽利用率 ├──MemoryL0.csv # L0A、L0B、L0C 的读写带宽利用率 ├──MemoryUB.csv # Vector 和 Scalar 到 UB 的读写带宽利用率 ├──OpBasicInfo.csv # 算子基本信息 ├──PipeUtilization.csv # 计算与搬运单元的耗时与占比 ├──ResourceConflictRatio.csv # UB bank 组、bank 冲突与资源冲突占比 └──visualize_data.bin # MindStudio Insight 展示文件例如通过cat ./OPPROF_*/PipeUtilization.csv即可查看 Task Duration 等关键指标定位算子性能瓶颈。八、支持产品与版本一览各入门样例对硬件与 CANN 版本有明确要求汇总如下样例支持产品CANN 版本要求00_quickstart / 01_add / 02_matrixbasic、advanced/ 03_fusion_operationAscend 950PR/950DT、Atlas A2/A3 训练与推理系列950PR/950DT ≥ CANN 9.1.0A2/A3 ≥ CANN 9.0.002_matrix/matmul_tensor_apiAscend 950PR/Ascend 950DT≥ CANN 9.1.004_reg_computeAscend 950PR/Ascend 950DT≥ CANN 9.1.0九、总结一条可复制的入门路径从 00_quickstart 的核函数直调起步到 01_add 掌握静态 Tensor 与 TQue/TPipe 两套向量编程范式再到 02_matrix 与 03_fusion_operation 理解 Cube 矩阵乘与跨单元融合最后在 04_reg_compute 接触寄存器级编程——01_simd_cpp_api目录以五个递进场景完整覆盖了 SIMD C API 的核心能力。每个样例都遵循统一的工程结构.asc内核源文件 CMakeLists.txt构建工程 数据生成/校验脚本与统一的编译选项CMAKE_ASC_RUN_MODE、CMAKE_ASC_ARCHITECTURES跑通任一样例即可举一反三。在此基础上可进一步深入仓库 1_Features 的硬件特性与 2_Performance 的性能演进故事完成从入门跑通到体系化调优的进阶。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
Ascend C SIMD C++ API 入门实战：cann-samples 核函数直调、向量计算、矩阵乘、融合与 RegBase 编程样例解析
Ascend C SIMD C API 入门实战cann-samples 核函数直调、向量计算、矩阵乘、融合与 RegBase 编程样例解析【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samplescann-samples 仓库的 01_simd_cpp_api 目录集中提供了基于 Ascend C SIMD C API 的入门样例覆盖核函数直调、向量计算、矩阵乘计算、融合计算以及 RegBase 向量编程五大基础场景。本文以该目录的官方概览文档为主线结合仓库内的样例源码、构建脚本与运行说明逐层拆解每个样例的工程结构、编程范式与编译运行流程帮助开发者在 Ascend 硬件上快速上手 SIMD C API并掌握加载-计算-存储三段式流水、多核并行切分、队列式内存管理等核心编程模式。一、样例目录定位从零开始的 SIMD C API 学习路径01_simd_cpp_api 是 cann-samples 中 SIMD单指令多数据C API 的入门级样例集合。与仓库内更高阶的性能故事如 2_Performance 下的各类性能调优实战不同本目录的定位是帮助开发者理解样例工程结构、编译运行流程和常见编程模式聚焦跑通第一个算子这一核心目标。其覆盖的五类基础场景正好构成一条由浅入深的学习主线核函数直调用内核调用符在 NPU 上启动核函数理解最小可运行工程向量计算以 Add 算子为例对比两种不同的内存与同步管理机制矩阵乘计算接触 Cube 单元的 Matmul 计算对比基础 API 与高级 API 两种实现融合计算将 Matmul 矩阵乘与 LeakyRelu 激活函数融合减少中间结果搬运RegBase 向量编程寄存器基RegBase编程范式下的向量计算与 VFVector Function融合。样例总览根据官方概览文档本目录包含的样例及其功能描述如下目录名称功能描述00_quickstartHelloWorld 核函数实现01_add基于不同编程模式的向量计算实现02_matrix矩阵乘计算03_fusion_operation融合计算04_reg_compute基于 RegBase 编程的向量计算从源码结构看各子目录均包含.asc内核源码文件、CMakeLists.txt构建文件与README.md说明文档向量计算类样例还附带scripts/gen_data.py与scripts/verify_result.py数据生成与校验脚本每个样例都遵循统一的实现 构建 运行 验证工程范式可直接作为新算子开发的脚手架。二、00_quickstartHelloWorld 核函数直调00_quickstart 是整套样例的起点通过 hello_world 样例演示了核函数在 NPU 侧运行验证的基础流程使用内核调用符启动核函数核函数内通过printf打印输出结果。2.1 最小内核工程结构Samples/0_Introduction/01_simd_cpp_api/00_quickstart/hello_world/ ├── CMakeLists.txt // 编译工程文件 ├── hello_world.asc // Ascend C 样例实现 调用样例 └── README.md // 样例说明文档其中 hello_world.asc 完整呈现了 Ascend C 内核开发的最小闭环#include utils/debug/asc_printf.h #include acl/acl.h __global__ __vector__ void hello_world() { printf(Hello World!!!\n); } int main(int argc, char const* argv[]) { aclrtSetDevice(0); // 运行管理资源申请。 aclrtStream stream nullptr; aclrtCreateStream(stream); hello_world8, 0, stream(); aclrtSynchronizeStream(stream); aclrtDestroyStream(stream); aclrtResetDevice(0); // 销毁运行资源。 return 0; }这段代码体现了 SIMD C API 内核直调的三个关键要素__global__ __vector__修饰符声明这是一个在 AIVAI Vector核上执行的向量内核函数printf输出能力来自 asc_printf.h 头文件8, 0, stream内核调用符第一个参数8指定使用 8 个核并行执行第二个参数0表示不使用局部内存参数第三个参数指定 ACL Stream源码位置ACL 运行时管理aclrtSetDevice/aclrtCreateStream申请设备与流资源aclrtSynchronizeStream等待内核执行完成aclrtDestroyStream/aclrtResetDevice释放资源构成完整的资源生命周期。2.2 构建系统与编译选项CMakeLists.txt 展示了 SIMD C API 工程的构建方式通过find_package(ASC REQUIRED)引入 Ascend 编译工具链以project(kernel_samples LANGUAGES ASC)声明工程语言再用add_executable(demo hello_world.asc)将.asc内核源文件编译为可执行程序并通过--npu-arch${CMAKE_ASC_ARCHITECTURES}编译选项指定 NPU 架构。本目录所有样例统一支持以下两个编译选项选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2/A3 训练与推理系列产品dav-3510对应 Ascend 950PR/Ascend 950DT2.3 编译、运行与结果验证在样例根目录执行以下步骤默认 npu 模式mkdir -p build cd build; cmake -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # 编译工程默认 npu 模式 ./demo若使用 CPU 调试或 NPU 仿真模式追加对应参数即可cmake -DCMAKE_ASC_RUN_MODEcpu -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # CPU 调试模式 cmake -DCMAKE_ASC_RUN_MODEsim -DCMAKE_ASC_ARCHITECTURESdav-2201 ..;make -j; # NPU 仿真模式注意切换编译模式前需清理 build 目录下的CMakeCache.txt缓存文件再重新执行 cmake。执行成功后输出 8 行打印每行对应一个 AIVAI VectorCore 的打印结果[AIV Block 0/8] Hello World!!! [AIV Block 1/8] Hello World!!! [AIV Block 2/8] Hello World!!! [AIV Block 3/8] Hello World!!! [AIV Block 4/8] Hello World!!! [AIV Block 5/8] Hello World!!! [AIV Block 6/8] Hello World!!! [AIV Block 7/8] Hello World!!![AIV Block X/8]表示当前为第 X 个核共 8 个核该输出证明核函数已在 8 个 AIV Core 上成功并行执行——这是理解后续多核并行计算样例每个核通过block_idx切分数据的基础。三、01_add向量计算的两种编程范式01_add 以 Add 算子z_i x_i y_i逐元素相加为载体验证 SIMD C API 的向量计算能力核心价值在于对比两种不同的内存与同步管理机制目录编程范式说明add静态 Tensor 编程使用LocalMemAllocator直接管理 UB 内存手动插入PipeBarrier同步add_tpipe_tqueTQue / TPipe 队列式编程通过EnQue/DeQue队列接口管理内存与流水同步两个样例的输入输出规格一致x、y为float类型、shape[8, 2048]、ND 布局输出z与输入同规格均以 8 核并行完成计算每核处理 2048 个元素blockLength 2048总计 8×204816384 个 float 元素。3.1 静态 Tensor 范式三段式流水线add 样例的计算逻辑遵循加载Load-计算Compute-存储Store三段式流水结构加载将输入数据 x、y 从 GMGlobal Memory片外全局内存容量大但访问慢搬运到 UBUnified BufferAI Core 片内向量计算专用缓存容量小但访问快经GlobalTensor与LocalTensor访问计算在 UB 上对xLocal与yLocal执行向量加法结果存入zLocal存储将计算结果从 UB 写回 GM。其核心实现涉及四个关键 API/变量DataCopyGM 与 UB 之间的数据搬运接口搬运方向由参数顺序决定PipeBarrierPIPE_ALL()流水线同步屏障确保数据搬运完成后才能执行后续操作避免读写冲突block_idx内置变量表示当前核的索引等价于GetBlockIdx()用于多核并行时的数据切分LocalMemAllocatorHardware::UBUB 片上内存分配器为向量计算分配连续的LocalTensor缓冲区。核心代码示意如下template uint32_t blockLength __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z) { AscendC::InitSocState(); // Global Tensor在 GM 上按 block_idx 偏移切分数据段实现多核并行 AscendC::GlobalTensorfloat xGm, yGm, zGm; xGm.SetGlobalBuffer(x block_idx * blockLength, blockLength); yGm.SetGlobalBuffer(y block_idx * blockLength, blockLength); zGm.SetGlobalBuffer(z block_idx * blockLength, blockLength); // Local Tensor在 UB 上分配计算缓冲区 AscendC::LocalMemAllocatorAscendC::Hardware::UB ubAllocator; AscendC::LocalTensorfloat xLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat yLocal ubAllocator.Allocfloat, blockLength(); AscendC::LocalTensorfloat zLocal ubAllocator.Allocfloat, blockLength(); // 阶段 1 加载GM - UB AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保加载完成后才能计算 // 阶段 2 计算z x y AscendC::Add(zLocal, xLocal, yLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保计算完成后才能存储 // 阶段 3 存储UB - GM AscendC::DataCopy(zGm, zLocal, blockLength); AscendC::PipeBarrierPIPE_ALL(); // 确保存储完成 }调用方式为numBlocks, 0, stream其中numBlocks8指定 8 核并行执行。整段实现以InitSocState()初始化 AI Core 硬件状态为起点SetGlobalBuffer(x block_idx * blockLength, ...)让每个核处理不同数据段三次PipeBarrierPIPE_ALL()分别保障加载→计算、计算→存储、存储完成的数据一致性。3.2 TQue/TPipe 范式队列式内存管理add_tpipe_tque 展示另一种编程思路——用 TPipe 和 TQue 统一管理内存分配与流水同步这是更接近生产级算子的写法。其处理流程为add_custom作为内核入口接收totalLength通过GetBlockNum()计算当前块需处理的数据长度用GetBlockIdx()计算当前核在 GM 中的起始位置用DataCopy将输入从 GM 搬入 UB并通过EnQue将输入LocalTensor放入输入队列通过DeQue从输入队列取出张量在 UB 上执行Add再通过EnQue将结果LocalTensor放入输出队列通过DeQue从输出队列取出结果用DataCopy写回当前核负责的 GM 数据段。该样例工程还包含完整的数据验证链路scripts/gen_data.py生成输入与真值数据scripts/verify_result.py校验输出与真值的一致性运行流程为mkdir -p build cd 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 # 校验输出结果与算法逻辑校验通过时输出test pass!。3.3 从 Add 到性能优化可优化方向分析官方文档还从实战角度给出了 Add 算子的四个可优化方向对理解后续性能样例极具启发序号优化方向当前实现问题预期收益1多核动态分配固定使用 8 核未按实际可用核数动态分配动态获取可用核数充分利用多核并行降低端到端时延2增大搬运粒度每次搬运 2048 个 float8KB粒度偏小增大单次搬运数据量、减少搬运次数摊薄启动开销提升带宽利用率3双缓冲流水并行加载、计算、存储严格串行MTE2/V/MTE3 硬件单元无法同时工作采用 Ping-Pong 双缓冲机制使三段并行、隐藏搬运时延4L2 Cache 旁路输入数据仅读取一次却默认经过 L2 Cache增加 Cache 污染对流式访问数据设置 L2 Cache 旁路减少无效 Cache 开销提升搬运效率四、02_matrix矩阵乘计算的三种 API 实现02_matrix 将场景从向量计算扩展到矩阵乘计算演示了 SIMD C API 面向 Cube 单元的三种实现路径目录实现方式支持产品matmul_basic_api基于静态 Tensor 编程范式实现矩阵乘950PR/950DT、Atlas A2/A3 训练与推理系列matmul_advanced_api基于 Matmul 高级 API 实现矩阵乘950PR/950DT、Atlas A2/A3 训练与推理系列matmul_tensor_api基于静态 Tensor API 编程范式实现矩阵乘950PR/950DT三个样例matmul_basic_api 与 matmul_advanced_api 均含scripts/gen_data.py、scripts/verify_result.py与 data_utils.h 数据工具头文件共享相同的数据生成与精度校验流程差异体现在内核侧 API 层次基础 API要求开发者显式管理 GlobalTensor/LocalTensor 与数据搬运Matmul 高级 API将切块、搬运与 Cube 计算封装为高层接口显著降低开发门槛静态 Tensor API则提供另一套静态张量编程范式。开发者可据此在同一计算目标下对比不同 API 的表达力与灵活性。五、03_fusion_operationCube 与 Vector 融合计算03_fusion_operation 引入自定义融合计算场景将 Matmul 矩阵乘与 LeakyRelu 激活函数融合为单算子避免中间结果落回 GM 带来的额外搬运开销目录功能描述支持产品matmul_leakyrelu_advanced_api基于高级 API 实现 Matmul 矩阵乘与 LeakyRelu 激活函数的融合计算950PR/950DT、Atlas A2/A3 训练与推理系列matmul_leakyrelu_basic_api基于基础 API 实现 Matmul 矩阵乘与 LeakyRelu 激活函数的融合计算950PR/950DT、Atlas A2/A3 训练与推理系列从实现角度看融合算子内部呈现典型的Cube矩阵乘 Vector激活函数协同流水Cube 单元完成矩阵乘后在片上直接衔接 Vector 单元的 LeakyRelu 计算两个样例分别验证了在基础 API 与高级 API 两种编程范式下完成跨单元融合的可行性是理解算子融合收益减少片内外数据往返、降低端到端时延的入门范例。六、04_reg_computeRegBase 向量编程04_reg_compute 介绍 SIMD C API 中RegBase寄存器基编程的向量计算能力目前仅支持 Ascend 950PR/Ascend 950DT目录功能描述add基于 RegBase 编程接口的 Add 示例gelu基于 LocalMemAllocator 与 RegBase/VF 融合的 GELU 计算示例RegBase 编程将数据与计算直接锚定在寄存器/寄存器组层面相比默认的内存基MemBase编程可获得更精细的数据布局与计算控制。其中 gelu 样例进一步将 RegBase 与 VFVector Function融合使用配合LocalMemAllocator完成片上内存分配展示了 RegBase 在非线性激活函数场景下的组合编程能力。对于希望深入寄存器级向量优化的开发者这是通往仓库内 Reg数据搬运场景选型指南 与 RegBase 性能故事如 softmax_regbase_story、kv_rms_norm_rope_cache_story的桥梁。七、贯穿全目录的调试与性能分析方法所有 SIMD C API 入门样例共享同一套调试与性能分析手段官方在 add 文档中给出了完整说明7.1 功能调试printf 与 DumpTensorAscendC::printf在内核侧需要输出日志的位置调用例如AscendC::printf(add blockIdx%d\n, AscendC::GetBlockIdx());适用于 CPU/NPU 域调试AscendC::DumpTensor用于 Dump 指定LocalTensor的内容并支持打印附加的自定义信息仅支持 uint32_t 类型例如在Add计算后调用AscendC::DumpTensor(zLocal, 1, 32);打印前 32 个元素。注意printf/DumpTensor 打印功能对算子实际运行性能有一定影响通常仅在调试阶段使用可通过设置ASCENDC_DUMP0关闭打印功能。7.2 性能调试msOpProf 单算子性能分析工具msOpProf 是单算子性能分析工具支持msopprof与msopprof simulator两种模式可对不同运行模式设备端/仿真和文件类型可执行程序/算子二进制.o文件进行性能数据采集与自动解析。设备端采集直接度量算子在 Ascend AI 处理器上的真实执行时间运行msopprof ./demo即可。命令执行后默认目录下生成OPPROF_{timestamp}_XXX文件夹其关键性能数据文件包括├──dump # 原始性能数据用户无需查看 ├──ArithmeticUtilization.csv # Cube/Vector 指令周期占比 ├──L2Cache.csv # L2 Cache 命中率影响 MTE2需合理规划数据搬运逻辑 ├──Memory.csv # UB、L1、主存的读写带宽利用率 ├──MemoryL0.csv # L0A、L0B、L0C 的读写带宽利用率 ├──MemoryUB.csv # Vector 和 Scalar 到 UB 的读写带宽利用率 ├──OpBasicInfo.csv # 算子基本信息 ├──PipeUtilization.csv # 计算与搬运单元的耗时与占比 ├──ResourceConflictRatio.csv # UB bank 组、bank 冲突与资源冲突占比 └──visualize_data.bin # MindStudio Insight 展示文件例如通过cat ./OPPROF_*/PipeUtilization.csv即可查看 Task Duration 等关键指标定位算子性能瓶颈。八、支持产品与版本一览各入门样例对硬件与 CANN 版本有明确要求汇总如下样例支持产品CANN 版本要求00_quickstart / 01_add / 02_matrixbasic、advanced/ 03_fusion_operationAscend 950PR/950DT、Atlas A2/A3 训练与推理系列950PR/950DT ≥ CANN 9.1.0A2/A3 ≥ CANN 9.0.002_matrix/matmul_tensor_apiAscend 950PR/Ascend 950DT≥ CANN 9.1.004_reg_computeAscend 950PR/Ascend 950DT≥ CANN 9.1.0九、总结一条可复制的入门路径从 00_quickstart 的核函数直调起步到 01_add 掌握静态 Tensor 与 TQue/TPipe 两套向量编程范式再到 02_matrix 与 03_fusion_operation 理解 Cube 矩阵乘与跨单元融合最后在 04_reg_compute 接触寄存器级编程——01_simd_cpp_api目录以五个递进场景完整覆盖了 SIMD C API 的核心能力。每个样例都遵循统一的工程结构.asc内核源文件 CMakeLists.txt构建工程 数据生成/校验脚本与统一的编译选项CMAKE_ASC_RUN_MODE、CMAKE_ASC_ARCHITECTURES跑通任一样例即可举一反三。在此基础上可进一步深入仓库 1_Features 的硬件特性与 2_Performance 的性能演进故事完成从入门跑通到体系化调优的进阶。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考