CANN Ascend C 入门实战基于 TPipe 与 TQue 的 Add 向量加法样例深度解析【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples本篇技术指南以 cann-samples 仓库中 add_tpipe_tque 样例 为主体系统讲解 Ascend C 中最基础的队列式TQue/TPipe编程范式从核函数如何按核切分 GM 数据、借助DataCopy搬入 UB、通过EnQue/DeQue完成阶段间数据交接到执行Add后写回 GM 的完整链路。读完本文你将掌握 TPipe/TQue 队列编程的标准模板、对应的 ACL 宿主侧调用方式、三种编译运行模式以及结果验证与调优的完整实操方法。样例定位入门系列中的“队列式编程”范本在 cann-samples 的 01_simd_cpp_api 入门目录中01_add小节集中展示了向量加法在不同编程模式下的实现方式add基于静态 Tensor编程LocalMemAllocator直接申请 UB 空间配合PipeBarrier做流水线同步add_tpipe_tque基于TPipe 和 TQue的队列式编程即本文主角。两者的计算逻辑完全一致z x y区别仅在于内存申请与同步方式静态 Tensor 方式显式调用LocalMemAllocator::Alloc并手动插入PipeBarrierPIPE_ALL而 TQue/TPipe 方式将 UB 缓冲交给TPipe::InitBuffer统一管理用EnQue/DeQue在“搬入—计算—搬出”各阶段之间传递LocalTensor由队列机制承担同步职责。理解这两个范本的差异是掌握 Ascend C 入门阶段最重要的分水岭。支持的产品与 CANN 软件版本产品CANN 软件版本Ascend 950PR/Ascend 950DT CANN 9.1.0Atlas A3 训练系列产品/Atlas A3 推理系列产品 CANN 9.0.0Atlas A2 训练系列产品/Atlas A2 推理系列产品 CANN 9.0.0从 CMakeLists.txt 可以看到样例默认的 NPU 架构为dav-2201对应 Atlas A2/A3 系列同时支持dav-3510对应 Ascend 950PR/Ascend 950DT编译时通过--npu-arch选项传递给 Bisheng 编译器。目录结构一览Samples/0_Introduction/01_simd_cpp_api/01_add/add_tpipe_tque ├── scripts │ ├── gen_data.py // 输入数据和真值数据生成脚本 │ └── verify_result.py // 验证输出数据和真值数据是否一致的验证脚本 ├── CMakeLists.txt // 编译工程文件 ├── data_utils.h // 数据读入写出函数 ├── add_tpipe_tque.asc // Ascend C 样例实现TPipe 和 TQue 管理内存和同步 调用样例 └── README.md // 样例说明文档其中 add_tpipe_tque.asc 是核心源码单个文件内同时包含 Device 侧核函数add_custom与 Host 侧调用main便于入门开发者通读完整链路。样例规格速览计算公式z x y即两组同形状数据的逐元素相加。输入输出张量shape数据类型格式x输入[8, 2048]floatNDy输入[8, 2048]floatNDz输出[8, 2048]floatND核函数名add_custom并行策略启动 8 个核并行计算每个核负责处理其中一段连续数据每核 2048 个元素。处理流程与队列编程范式样例的整体处理流程如下add_custom作为核入口接收totalLength全局输入长度在核函数内通过GetBlockNum()计算当前 block 应处理的数据长度blockLength通过GetBlockIdx()计算当前核在 GM 中对应的数据起点使用DataCopy把输入数据从 GM 搬到 UB并通过EnQue将输入LocalTensor放入输入队列通过DeQue从输入队列取出输入张量在 UB 中执行Add再通过EnQue将结果LocalTensor放入输出队列通过DeQue从输出队列取出结果使用DataCopy写回当前核负责的 GM 分片。这里的核心概念是TPipe与TQueTPipe负责管理 UB 内存缓冲与流水线同步的“管道”对象通过InitBuffer为各队列申请与数据长度匹配的 UB 空间TQue基于特定位置TPosition的队列本样例使用VECIN向量输入与VECOUT向量输出两种队列分别承载“待计算”与“待写回”的张量EnQue/DeQueEnQue将已搬到 UB 的LocalTensor入队DeQue在后续阶段从队列取出张量继续处理队列机制隐式承担了阶段间的数据依赖同步。一句话概括TPipe/TQue 把“内存申请、数据搬运、同步”等样板代码封装成队列操作让开发者把精力集中在计算逻辑本身。源码级解析核函数 add_custom完整实现见 add_tpipe_tque.asc下面分段拆解。1. 核函数声明与对象创建__global__ __vector__ void add_custom(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t* z, uint32_t totalLength) { AscendC::TPipe pipe; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueueX; AscendC::TQueAscendC::TPosition::VECIN, 1 inQueueY; AscendC::TQueAscendC::TPosition::VECOUT, 1 outQueueZ; AscendC::GlobalTensorfloat xGm; AscendC::GlobalTensorfloat yGm; AscendC::GlobalTensorfloat zGm;__global__ __vector__声明这是一个向量指令域的内核函数形参直接使用__gm__ uint8_t*裸地址模板参数TPosition::VECIN, 1中的1表示队列深度buffer num本样例每个队列只申请 1 块缓冲即“单缓冲、串行三段式”的最简形态GlobalTensorfloat用于描述 GM 上的张量视图。2. 按核切分数据区间uint32_t blockLength totalLength / AscendC::GetBlockNum(); xGm.SetGlobalBuffer((__gm__ float*)x blockLength * AscendC::GetBlockIdx(), blockLength); yGm.SetGlobalBuffer((__gm__ float*)y blockLength * AscendC::GetBlockIdx(), blockLength); zGm.SetGlobalBuffer((__gm__ float*)z blockLength * AscendC::GetBlockIdx(), blockLength);GetBlockNum()返回本次启动的核数本例为 8blockLength 16384 / 8 2048GetBlockIdx()返回当前核的编号blockLength * GetBlockIdx()即当前核在 GM 中的起始元素偏移确保各核访问互不重叠的数据分片代码中的除法假定totalLength能被核数整除——这是该样例成立的前提实际算子需额外处理尾块。3. 申请 UB 缓冲并入队输入pipe.InitBuffer(inQueueX, 1, blockLength * sizeof(float)); pipe.InitBuffer(inQueueY, 1, blockLength * sizeof(float)); pipe.InitBuffer(outQueueZ, 1, blockLength * sizeof(float)); AscendC::LocalTensorfloat xLocal inQueueX.AllocTensorfloat(); AscendC::LocalTensorfloat yLocal inQueueY.AllocTensorfloat(); AscendC::DataCopy(xLocal, xGm, blockLength); AscendC::DataCopy(yLocal, yGm, blockLength); inQueueX.EnQue(xLocal); inQueueY.EnQue(yLocal);InitBuffer以“字节数”为单位为每个队列申请 UB 缓冲AllocTensor从队列缓冲中取出一块可用的LocalTensorDataCopy将其从 GM 搬入 UB随后EnQue把“已就绪”的张量推入队列——入队即通知下游阶段数据可用。4. 出队、计算、再入队xLocal inQueueX.DeQuefloat(); yLocal inQueueY.DeQuefloat(); AscendC::LocalTensorfloat zLocal outQueueZ.AllocTensorfloat(); AscendC::Add(zLocal, xLocal, yLocal, blockLength); outQueueZ.EnQuefloat(zLocal); inQueueX.FreeTensor(xLocal); inQueueY.FreeTensor(yLocal);DeQue从输入队列取出张量此时可安全执行 UB 内的向量加Add(zLocal, xLocal, yLocal, blockLength)计算结果EnQue到输出队列等待写回阶段取用输入张量处理完毕后调用FreeTensor释放归还给 TPipe 复用。5. 结果写回 GMzLocal outQueueZ.DeQuefloat(); AscendC::DataCopy(zGm, zLocal, blockLength); outQueueZ.FreeTensor(zLocal);从输出队列取出结果DataCopy写回当前核负责的 GM 分片最后FreeTensor收尾。整个核函数没有出现任何显式的PipeBarrier——这正是 TQue/TPipe 的价值阶段间的同步由队列机制自动完成。对比同一目录下基于静态 Tensor 的 add.asc后者需要手动调用AscendC::PipeBarrierPIPE_ALL()来保证DataCopy与Add之间的数据就绪关系两种风格在源码层面形成了直观对照。宿主侧调用实现ACL 完整流程核函数之外add_tpipe_tque.asc 的main函数展示了完整的 Host 侧链路初始化与资源创建aclInit(nullptr)→aclrtSetDevice(deviceId)→aclrtCreateStream(stream)内存分配aclrtMallocHost分配 Host 侧内存aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分配 Device 侧内存数据准备ReadFile从./input/input_x.bin、./input/input_y.bin读入输入aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE)拷贝到 Device内核调用add_customnumBlocks, 0, stream(xDevice, yDevice, zDevice, totalLength);内核调用符中依次传入 block 数8、l2 参数0和 stream运行时参数传入 Device 侧 x、y、z 地址与总数据长度totalLength 5.同步与回拷aclrtSynchronizeStream(stream)等待内核完成aclrtMemcpy(..., ACL_MEMCPY_DEVICE_TO_HOST)回拷结果WriteFile写出./output/output.bin 6.资源释放依次aclrtFree、aclrtFreeHost、aclrtDestroyStream、aclrtResetDevice、aclFinalize。数据读写函数封装在 data_utils.h 中ReadFile基于statstd::ifstream做二进制读取含文件存在性、类型、空文件、缓冲区越界检查WriteFile基于open/write写出为多个入门样例所共用。数据生成与结果验证脚本数据生成gen_data.py 使用np.random.uniform(1, 10, [8, 2048])生成两个 float32 输入真值golden (input_x input_y).astype(np.float32)分别落盘为input/input_x.bin、input/input_y.bin与output/golden.bin。结果验证verify_result.py 采用np.isclose做逐元素比对容差参数为相对容差RELATIVE_TOL 1e-4绝对容差ABSOLUTE_TOL 1e-5允许的错误比例ERROR_TOL 1e-4即错误元素占比不超过 0.01% 即判定通过。脚本会打印每个差异元素的索引、期望值、实际值与相对偏差最多 100 个最终输出整体error ratio并返回test pass!或报错退出。编译与运行1. 配置环境变量根据当前环境上 CANN 开发套件包的安装方式参见仓库内 third_party/asc-devkit 子模块文档配置环境变量source ${install_path}/cann/set_env.sh说明${install_path}为 CANN 包安装目录未指定安装目录时默认安装至/usr/local/Ascend下。仓库的 cmake/ascend.cmake 会优先读取ASCEND_HOME_PATH环境变量定位工具链未设置时会尝试/usr/local/Ascend/ascend-toolkit/latest等默认路径。2. 样例执行NPU 模式默认在样例根目录执行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 # 验证输出结果是否正确确认算法逻辑正确其中 CMakeLists.txt 通过find_package(ASC REQUIRED)引入 Ascend 工具链add_executable(demo add_tpipe_tque.asc)直接以.asc为源构建可执行程序并用--npu-arch${CMAKE_ASC_ARCHITECTURES}指定 NPU 架构。3. CPU 调试与 NPU 仿真模式使用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。4. 编译选项说明选项可选值说明CMAKE_ASC_RUN_MODEnpu默认、cpu、sim运行模式NPU 运行、CPU 调试、NPU 仿真CMAKE_ASC_ARCHITECTURESdav-2201默认、dav-3510NPU 架构dav-2201对应 Atlas A2 训练系列产品/Atlas A2 推理系列产品和 Atlas A3 训练系列产品/Atlas A3 推理系列产品dav-3510对应 Ascend 950PR/Ascend 950DT5. 执行结果执行完毕后若精度比对通过终端输出test pass!小结与进阶路径本文围绕 cann-samples 中的add_tpipe_tque样例完整拆解了 TQue/TPipe 队列式编程的“切分—搬入—入队—出队—计算—入队—出队—写回”标准流水并结合 add_tpipe_tque.asc 源码、CMakeLists.txt 编译配置与 scripts 验证脚本给出了可直接复跑的完整闭环。以此为起点可以继续探索仓库内同目录的 静态 Tensor 实现 对比两种编程范式再进阶到 矩阵乘样例、融合算子样例 以及 RegBase 向量编程逐步建立 Ascend C 高性能算子开发的全景知识。【免费下载链接】cann-samplesCANN高性能实战演进样例与体系化调优知识库项目地址: https://gitcode.com/cann/cann-samples创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考