企业官网建设流程全解析

Ascend C Reg API 教程与高性能实现

第一部分：基础概念

1.1 什么是 Reg API

Reg API（寄存器级向量计算 API）是 Ascend C 在 CANN v9.0.0 引入的面向 SIMD Vector 寄存器的编程接口，对应硬件 SIMD Vector Core 内部的VF Reg（Vector Function Register）文件。所有 Reg API 都位于AscendC::Reg命名空间，使用__simd_callee__ inline修饰。

与传统基础 API（AscendC::Add、AscendC::Mul等 LocalTensor 接口）相比：

1.2 何时使用 Reg API

推荐使用 Reg API 的场景：

1. 多步融合计算：同一个数据要经过多个算子（如Cast → Mul → Add → Cast），中间结果可留在寄存器中避免反复 Load/Store UB
2. 计算密集型算子：如Exp/Ln/Sqrt/Div等指令周期较长，减少 UB 往返收益大
3. RegTrait 双宽场景：64 位整数运算（int64_t）需要RegTraitNumTwo把 2 个 VL 拼成 2×VL 的逻辑寄存器
4. 极致性能调优：基础 API 已达瓶颈，需要软件掌控指令调度
5. 非连续数据访问：Gather/Scatter、Squeeze、Block Strided Load 等 MemBase 难以高效表达的模式

不推荐使用 Reg API 的场景：

1. 简单单步算子（如纯Add/Sub），基础 API 已足够快
2. 团队不熟悉寄存器编程模型，调试成本高
3. 数据搬运本身就是瓶颈（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, 950PR=256KB │ (LocalMem) │ bank=16group×3bank×4KB, vec一拍访问256B │ __ubuf__ └──────────────┬──────────────────────────────┘ │ LoadAlign / StoreAlign (vld/vst) ▼ ┌─────────────────────────────────────────────┐ VF Reg │ Vector Function Register file │ RegTensor<T> │ 单个寄存器宽度 = 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 Function，Reg API 的唯一执行环境） │ ├─ 可调用: __simd_callee__ 函数（所有 AscendC::Reg::* 都是 __simd_callee__） │ ▼ __simd_callee__ 函数（叶子节点，仅能调用其他 __simd_callee__）

关键规则：

• Reg API（AscendC::Reg::*）全部标注__simd_callee__ inline
• __simd_vf__标记的函数只能通过asc_vf_call<VFFunc>(args...)调用
• __simd_callee__只能被__simd_vf__、其他__simd_callee__或constexpr aicore调用
• Reg API 不能在__aicore__函数中直接调用，必须经asc_vf_call进入 VF 上下文

合法调用模板：

// 1. 定义 SIMD VF 函数template<typenameT>__simd_vf__inlinevoidAddVF(__ubuf__ T*dst,__ubuf__ T*src0,__ubuf__ T*src1,uint16_trepeatTimes,uint32_toneRepeatSize){AscendC::Reg::RegTensor<T>r0,r1,r2;AscendC::Reg::MaskReg mask=AscendC::Reg::CreateMask<T>();for(uint16_ti=0;i<repeatTimes;++i){AscendC::Reg::LoadAlign(r0,src0+i*oneRepeatSize);AscendC::Reg::LoadAlign(r1,src1+i*oneRepeatSize);AscendC::Reg::Add(r2,r0,r1,mask);// __simd_callee__ 调用AscendC::Reg::StoreAlign(dst+i*oneRepeatSize,r2,mask);}}// 2. 在 __aicore__ 中调用__aicore__inlinevoidProcess(){// ... DataCopy / SetFlag / WaitFlag 等 MTE 操作 ...asc_vf_call<AddVF<T>>(dstAddr,src0Addr,src1Addr,repeatTimes,oneRepeatSize);// ... 后续 MTE3 ...}

第二部分：Reg API 数据类型与寄存器

2.1 RegTensor —— 向量数据寄存器

RegTensor<T>是 Reg API 的核心数据载体，定义在kernel_reg_compute_struct_intf.h：

structRegTrait{intREG_NUM=1;};constexprRegTrait RegTraitNumOne={1};// 1 个 VL 宽寄存器（默认）constexprRegTrait RegTraitNumTwo={2};// 2 个 VL 宽寄存器（int64 等需要）template<typenameT,constRegTrait&regTrait=RegTraitNumOne>structRegTensor{usingActualT=T;usingRegType=typenameTypeGet<T>::T;RegType reg[regTrait.REG_NUM];// 物理寄存器数组staticconstexprintREG_NUM=regTrait.REG_NUM;};

容量：RegTensor<T>一次可装载VL / sizeof(T)个元素：

• float/half/int32_t：VL=256B → 64 个 float / 128 个 half / 64 个 int32
• int64_t：需用RegTraitNumTwo拼接 2 个 VL → 64 个 int64（逻辑上 512B 宽寄存器）

使用要点：

• RegTensor<T>应声明在__simd_vf__函数内部，作为栈上变量
• 寄存器生命周期由硬件 rename 表管理，软件声明只是逻辑寄存器
• dstReg和srcReg可以是同一个变量（如Add(r, r, r, mask)），硬件会正确处理 RAW 依赖

2.2 MaskReg —— 谓词掩码寄存器

MaskReg = vector_bool，宽度为VL/8（950PR 都是 32B = 256 bit）。每一位对应RegTensor中的一个 lane（最小粒度为字节级 mask 的部分场景除外）。

创建 Mask 的两种方式：

// 1. CreateMask: 编译期固定模式AscendC::Reg::MaskReg maskAll=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::ALL>();// 全 1AscendC::Reg::MaskReg maskVL1=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::VL1>();// 仅最低 1 laneAscendC::Reg::MaskReg maskVL64=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::VL64>();// 低 64 laneAscendC::Reg::MaskReg maskM3=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::M3>();// 每 4 lane 取 3AscendC::Reg::MaskReg maskH=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::H>();// 高半AscendC::Reg::MaskReg maskQ=AscendC::Reg::CreateMask<T,AscendC::Reg::MaskPattern::Q>();// 1/4// 2. UpdateMask: 运行期根据剩余 count 生成uint32_tcount=200;// 需要处理的元素数AscendC::Reg::MaskReg mask=AscendC::Reg::UpdateMask<T>(count);// 当 count >= VL/sizeof(T) 时返回 ALL mask// 当 count < VL/sizeof(T) 时返回低 count 位为 1 的 mask

MaskPattern 枚举全集（来自kernel_reg_compute_utils.h）：
ALL, VL1, VL2, VL3, VL4, VL8, VL16, VL32, VL64, VL128, M3, M4, H, Q, ALLF

MaskMergeMode —— mask 未命中位的处理策略：

典型场景：处理非对齐尾段时，最后一拍 repeat 用UpdateMask(remainder)生成部分 mask，配合MERGING模式避免越界写 0 污染已写入数据。详见样例mergemode/mergemode.asc：

AscendC::Reg::Duplicate(dstReg,(T)2);// 先填默认值mask=AscendC::Reg::UpdateMask<T>(count);// 生成部分 maskAscendC::Reg::LoadAlign<T,AscendC::Reg::PostLiteral::POST_MODE_UPDATE>(src0Reg,src0Addr,oneRepeatSize);AscendC::Reg::Max<T,AscendC::Reg::MaskMergeMode::MERGING>(dstReg,src0Reg,src1Reg,mask);AscendC::Reg::StoreAlign(dstAddr+i*oneRepeatSize,dstReg,allMask);

2.3 AddrReg —— 地址偏移寄存器

AddrReg = vector_address，存储 (index, stride) 二元组，配合LoadAlign/StoreAlign的 AddrReg 重载，可在循环中自动累加地址偏移，省去软件计算addr + i * size的开销。

// 4 个 (index, stride) 对的 CreateAddrReg 重载AddrRegCreateAddrReg<T>(uint16_tindex0,uint32_tstride0);AddrRegCreateAddrReg<T>(uint16_ti0,uint32_ts0,uint16_ti1,uint32_ts1);AddrRegCreateAddrReg<T>(uint16_ti0,uint32_ts0,uint16_ti1,uint32_ts1,uint16_ti2,uint32_ts2);AddrRegCreateAddrReg<T>(uint16_ti0,uint32_ts0,...,uint16_ti3,uint32_ts3);// 使用AscendC::Reg::AddrReg areg;for(uint16_ti=0;i<repeatTimes;++i){areg=AscendC::Reg::CreateAddrReg<T>(i,oneRepeatSize);// index=i, stride=oneRepeatSizeAscendC::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 数据；写回侧用StoreUnAlign+StoreUnAlignPost。

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 Register）：

AscendC::Reg::ClearSpr<AscendC::SpecialPurposeReg::AR>();// 清地址寄存器// 详见 docs/api/SIMD-API/基础API/Reg矢量计算/系统变量访问/