Ascend C Reg API 教程与高性能实现第一部分基础概念1.1 什么是 Reg APIReg API寄存器级向量计算 API是 Ascend C 在 CANN v9.0.0 引入的面向 SIMD Vector 寄存器的编程接口对应硬件 SIMD Vector Core 内部的VF RegVector Function Register文件。所有 Reg API 都位于AscendC::Reg命名空间使用__simd_callee__ inline修饰。与传统基础 APIAscendC::Add、AscendC::Mul等 LocalTensor 接口相比维度基础 APIMemBaseReg APIRegBase数据载体LocalTensorTUB 内存RegTensorTVF 寄存器单次处理粒度任意长度硬件内部 tiling一个VLVector Length950PR256B软件显式分块不需要需要循环每次处理 VL/sizeof(T) 个元素数据搬运DataCopyMTE2/MTE3LoadAlign/StoreAlign直接 UB↔VF Reg控制流__aicore__内联必须用__simd_vf__函数 asc_vf_call调用Mask 控制通过 count 参数自动显式MaskReg寄存器寄存器复用硬件决定软件可控dstsrc 时直接复用性能上限受 MTE 流水限制减少 UB 往返提升指令并发编程复杂度简单较高需理解寄存器/调用层级1.2 何时使用 Reg API推荐使用 Reg API 的场景多步融合计算同一个数据要经过多个算子如Cast → Mul → Add → Cast中间结果可留在寄存器中避免反复 Load/Store UB计算密集型算子如Exp/Ln/Sqrt/Div等指令周期较长减少 UB 往返收益大RegTrait 双宽场景64 位整数运算int64_t需要RegTraitNumTwo把 2 个 VL 拼成 2×VL 的逻辑寄存器极致性能调优基础 API 已达瓶颈需要软件掌控指令调度非连续数据访问Gather/Scatter、Squeeze、Block Strided Load 等 MemBase 难以高效表达的模式不推荐使用 Reg API 的场景简单单步算子如纯Add/Sub基础 API 已足够快团队不熟悉寄存器编程模型调试成本高数据搬运本身就是瓶颈mte2_ratio 90%的算子Reg API 收益有限来自实测数据Ascend 950 系列 Add 算子 RegBase 相对 MemBase 的主要收益来自减少冗余 Load/Store、寄存器复用、提升无依赖并发指令比例但当前 Add 样例核心计算基本是单条 Add计算链短、可融合/双发空间有限难以充分释放 RegBase 优势。这说明 Reg API 在多步融合场景才能体现价值。1.3 存储层级与编程模型┌─────────────────────────────────────────────┐ GM (HBM) │ ~1.6-1.8 TB/s │ (Global) └──────────────┬──────────────────────────────┘ │ DataCopy (MTE2/MTE3) ▼ ┌─────────────────────────────────────────────┐ UB │ Unified Buffer, 950PR256KB │ (LocalMem) │ bank16group×3bank×4KB, vec一拍访问256B │ __ubuf__ └──────────────┬──────────────────────────────┘ │ LoadAlign / StoreAlign (vld/vst) ▼ ┌─────────────────────────────────────────────┐ VF Reg │ Vector Function Register file │ RegTensorT │ 单个寄存器宽度 VL 256B │ │ RegTraitNumOne: 1 reg; NumTwo: 2 reg (int64)│ │ MaskReg (vector_bool): 32B (256 bits) │ │ AddrReg (vector_address): 地址偏移寄存器 │ │ UnalignRegForLoad/Store (vector_align): │ │ 非对齐访存辅助寄存器 │ └──────────────┬──────────────────────────────┘ │ ▼ AscendC::Reg::Add / Mul / Cast / ... VALU 算术逻辑单元执行1.4 调用层级Reg API 的调用链有严格层级违反会编译报错__global__ __aicore__ kernel (Host 入口) │ ├─ 可调用: __aicore__ 函数、constexpr aicore │ ▼ __aicore__ 函数算子主流程 │ ├─ 可调用: __simd_vf__ 函数 (通过 asc_vf_call)、__aicore__ 函数 │ ▼ __simd_vf__ 函数SIMD Vector FunctionReg API 的唯一执行环境 │ ├─ 可调用: __simd_callee__ 函数所有 AscendC::Reg::* 都是 __simd_callee__ │ ▼ __simd_callee__ 函数叶子节点仅能调用其他 __simd_callee__关键规则Reg APIAscendC::Reg::*全部标注__simd_callee__ inline__simd_vf__标记的函数只能通过asc_vf_callVFFunc(args...)调用__simd_callee__只能被__simd_vf__、其他__simd_callee__或constexpr aicore调用Reg API 不能在__aicore__函数中直接调用必须经asc_vf_call进入 VF 上下文合法调用模板// 1. 定义 SIMD VF 函数templatetypenameT__simd_vf__inlinevoidAddVF(__ubuf__ T*dst,__ubuf__ T*src0,__ubuf__ T*src1,uint16_trepeatTimes,uint32_toneRepeatSize){AscendC::Reg::RegTensorTr0,r1,r2;AscendC::Reg::MaskReg maskAscendC::Reg::CreateMaskT();for(uint16_ti0;irepeatTimes;i){AscendC::Reg::LoadAlign(r0,src0i*oneRepeatSize);AscendC::Reg::LoadAlign(r1,src1i*oneRepeatSize);AscendC::Reg::Add(r2,r0,r1,mask);// __simd_callee__ 调用AscendC::Reg::StoreAlign(dsti*oneRepeatSize,r2,mask);}}// 2. 在 __aicore__ 中调用__aicore__inlinevoidProcess(){// ... DataCopy / SetFlag / WaitFlag 等 MTE 操作 ...asc_vf_callAddVFT(dstAddr,src0Addr,src1Addr,repeatTimes,oneRepeatSize);// ... 后续 MTE3 ...}第二部分Reg API 数据类型与寄存器2.1 RegTensor —— 向量数据寄存器RegTensorT是 Reg API 的核心数据载体定义在kernel_reg_compute_struct_intf.hstructRegTrait{intREG_NUM1;};constexprRegTrait RegTraitNumOne{1};// 1 个 VL 宽寄存器默认constexprRegTrait RegTraitNumTwo{2};// 2 个 VL 宽寄存器int64 等需要templatetypenameT,constRegTraitregTraitRegTraitNumOnestructRegTensor{usingActualTT;usingRegTypetypenameTypeGetT::T;RegType reg[regTrait.REG_NUM];// 物理寄存器数组staticconstexprintREG_NUMregTrait.REG_NUM;};容量RegTensorT一次可装载VL / sizeof(T)个元素float/half/int32_tVL256B → 64 个 float / 128 个 half / 64 个 int32int64_t需用RegTraitNumTwo拼接 2 个 VL → 64 个 int64逻辑上 512B 宽寄存器使用要点RegTensorT应声明在__simd_vf__函数内部作为栈上变量寄存器生命周期由硬件 rename 表管理软件声明只是逻辑寄存器dstReg和srcReg可以是同一个变量如Add(r, r, r, mask)硬件会正确处理 RAW 依赖2.2 MaskReg —— 谓词掩码寄存器MaskReg vector_bool宽度为VL/8950PR 都是 32B 256 bit。每一位对应RegTensor中的一个 lane最小粒度为字节级 mask 的部分场景除外。创建 Mask 的两种方式// 1. CreateMask: 编译期固定模式AscendC::Reg::MaskReg maskAllAscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::ALL();// 全 1AscendC::Reg::MaskReg maskVL1AscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::VL1();// 仅最低 1 laneAscendC::Reg::MaskReg maskVL64AscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::VL64();// 低 64 laneAscendC::Reg::MaskReg maskM3AscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::M3();// 每 4 lane 取 3AscendC::Reg::MaskReg maskHAscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::H();// 高半AscendC::Reg::MaskReg maskQAscendC::Reg::CreateMaskT,AscendC::Reg::MaskPattern::Q();// 1/4// 2. UpdateMask: 运行期根据剩余 count 生成uint32_tcount200;// 需要处理的元素数AscendC::Reg::MaskReg maskAscendC::Reg::UpdateMaskT(count);// 当 count VL/sizeof(T) 时返回 ALL mask// 当 count VL/sizeof(T) 时返回低 count 位为 1 的 maskMaskPattern 枚举全集来自kernel_reg_compute_utils.hALL, VL1, VL2, VL3, VL4, VL8, VL16, VL32, VL64, VL128, M3, M4, H, Q, ALLFMaskMergeMode —— mask 未命中位的处理策略模式行为适用场景MaskMergeMode::ZEROING默认mask0 的 lane dst 写 0第一次写入、不需要保留旧值MaskMergeMode::MERGINGmask0 的 lane dst 保留原值多次部分写入累加、保留之前结果典型场景处理非对齐尾段时最后一拍 repeat 用UpdateMask(remainder)生成部分 mask配合MERGING模式避免越界写 0 污染已写入数据。详见样例mergemode/mergemode.ascAscendC::Reg::Duplicate(dstReg,(T)2);// 先填默认值maskAscendC::Reg::UpdateMaskT(count);// 生成部分 maskAscendC::Reg::LoadAlignT,AscendC::Reg::PostLiteral::POST_MODE_UPDATE(src0Reg,src0Addr,oneRepeatSize);AscendC::Reg::MaxT,AscendC::Reg::MaskMergeMode::MERGING(dstReg,src0Reg,src1Reg,mask);AscendC::Reg::StoreAlign(dstAddri*oneRepeatSize,dstReg,allMask);2.3 AddrReg —— 地址偏移寄存器AddrReg vector_address存储 (index, stride) 二元组配合LoadAlign/StoreAlign的 AddrReg 重载可在循环中自动累加地址偏移省去软件计算addr i * size的开销。// 4 个 (index, stride) 对的 CreateAddrReg 重载AddrRegCreateAddrRegT(uint16_tindex0,uint32_tstride0);AddrRegCreateAddrRegT(uint16_ti0,uint32_ts0,uint16_ti1,uint32_ts1);AddrRegCreateAddrRegT(uint16_ti0,uint32_ts0,uint16_ti1,uint32_ts1,uint16_ti2,uint32_ts2);AddrRegCreateAddrRegT(uint16_ti0,uint32_ts0,...,uint16_ti3,uint32_ts3);// 使用AscendC::Reg::AddrReg areg;for(uint16_ti0;irepeatTimes;i){aregAscendC::Reg::CreateAddrRegT(i,oneRepeatSize);// indexi, strideoneRepeatSizeAscendC::Reg::LoadAlign(xReg,xAddr,areg);// 等价于 LoadAlign(xReg, xAddr i * oneRepeatSize)AscendC::Reg::StoreAlign(zAddr,zReg,areg,mask);}适用场景循环中地址按固定 stride 递增且 stride 较大需要地址计算外提时。多数情况下addr i * oneRepeatSize已足够AddrReg 在 post-update 模式或跨多段数据交错访问时更优。2.4 UnalignRegForLoad / UnalignRegForStore —— 非对齐辅助寄存器UnalignRegForLoad UnalignRegForStore vector_align用于跨 32B/256B 边界的非对齐数据搬移。非对齐搬移分两步先用LoadUnAlignPre把跨边界数据缓存到 ureg再用LoadUnAlign拼接成完整 VL 数据写回侧用StoreUnAlignStoreUnAlignPost。AscendC::Reg::UnalignRegForLoad ureg0;AscendC::Reg::UnalignRegForStore ureg1;AscendC::Reg::LoadUnAlignPre(ureg0,srcAddr,areg);// 预取跨边界部分AscendC::Reg::LoadUnAlign(srcReg,ureg0,srcAddr,areg,0);// 完整加载AscendC::Reg::StoreUnAlign(dstAddr,srcReg,ureg1,areg);// 非对齐写缓存尾部// 循环结束后AscendC::Reg::StoreUnAlignPost(dstAddr,ureg1,areg);// 刷出最后残留适用场景输出数据长度不是 VL/sizeof(T) 的整数倍且不能用MaskReg简单处理的非对齐连续场景如 Squeeze 压缩输出。2.5 特殊寄存器访问GetSpr / ClearSpr少数场景需要读写特殊功能寄存器System Special Purpose RegisterAscendC::Reg::ClearSprAscendC::SpecialPurposeReg::AR();// 清地址寄存器// 详见 docs/api/SIMD-API/基础API/Reg矢量计算/系统变量访问/