REGBASE 是昇腾 950PR/950DT 矢量单元（AIV）新增的基于寄存器的编程模型，

开发者通过 __simd_vf__ 标记的 Vector Function 和 AscendC::Reg 命名空间下的 Reg API，直接操作 AIV 核内新加的一级 SIMD Register File（VF Reg），把一段连续矢量计算完整保留在寄存器里，只在进入和退出时与 Unified Buffer（UB）交互一次。

设计背景：MemBase 编程的三大瓶颈

在以往的 MemBase 矢量编程里，矢量计算单元的源操作数和目的操作数都直接落在 Local Memory（UB）上，每条矢量指令的执行过程都是"从 UB 读 → 在执行单元算 → 写回 UB"。当算子由多个矢量计算串联（如 dst = (a + b) * c + d）时，每一步的中间结果都必须先写回 UB，再被下一条指令重新读出，由此产生三个典型瓶颈：

• UB 读写带宽被反复占用：中间结果在 UB 上往返，访存压力随矢量计算条数线性放大；
• UB Bank 冲突概率显著提升：对 UB 的反复读写极大增加了 Bank 冲突；
• 关键路径被访存拉长：当计算指令本身执行时间很短，整体耗时被 UB 读写时间占满。
  REGBASE 的定位就是突破这个"每步落地 UB"的性能天花板——硬件在矢量执行单元前新增了一级 SIMD Register File，软件开放面向这一级的编程接口，让开发者自主控制数据搬运和计算调度，将中间结果留在寄存器中连续消费。

硬件架构与数据通路

AIV 是 AI Core 内负责矢量计算的核，参与 Reg 矢量计算的硬件单元包括三部分：

• Reg 矢量执行单元：从寄存器读数据，完成计算后写回寄存器；
• DMA 单元：负责寄存器和 UB 之间的数据搬运（LoadAlign / StoreAlign）；
• Aux Scalar：处理地址计算等标量辅助工作。
  三者实际执行时都归属 PIPE_V 流水，无数据依赖时可同时发射、并行执行；有寄存器依赖时由硬件按指令顺序保证正确性，但跨寄存器对同一 UB 区域的读写需要开发者显式插入同步（搬运单元与计算单元之间没有自动顺序约束）。

内存层级与数据通路

REGBASE 的内存层级自外向内依次为 GM（HBM）→ UB（950PR 为 256 KB）→ VF Reg，且有一条硬性规则：寄存器不支持直接从 GM 加载数据或直接将数据写出 GM，数据流向必须经过 UB 中转，即 GM → UB → Reg → 计算 → Reg → UB → GM。UB 和寄存器均为每个 AIV 核内独享的存储空间，核内所有 VF Reg 共享同一个 UB 资源。
关于你提到的"矢量记忆"：昇腾的内存体系中并没有一个独立的"矢量记忆"模块，通常指的是 VF Reg（Vector Function Register File，矢量寄存器堆）和 UB（Unified Buffer，统一缓冲）这一对紧耦合的存储层级——VF Reg 是矢量执行单元的私有寄存器堆，UB 是片上本地存储，两者共同构成矢量计算的"近端记忆"。

寄存器分类

寄存器作为整体使用，不能按索引或偏移访问其中的某个元素或 bit。按功能主要分为以下几类：

寄存器类型	宽度	作用	与 UB 交换
矢量计算寄存器	VL（950PR = 256 B）	参与矢量计算的主要载体	支持（经 DMA）
掩码寄存器	VL/8	控制参与计算的有效元素	支持（经 DMA）
地址寄存器	32 bit	存储 UB 地址偏移，辅助搬运	不支持
搬入非对齐寄存器	DataBlock	临时存放非对齐源地址的对齐数据	支持（经 DMA）
其中 VL（Vector Length）是单个矢量计算寄存器的宽度，也是 REGBASE 每条指令一次能处理的最大数据长度，也是软件分块循环的步长基准。

编程模型与调用层级

REGBASE 的核心抽象是 Vector Function，使用 __simd_vf__ 标记，被发送到硬件中的矢量运算单元执行；函数内部通过 Reg 矢量计算 API（位于 AscendC::Reg 命名空间，以 simd_callee inline 修饰）完成计算操作。

调用层级对比

• Membase：核函数 → 基础 API（LocalTensor，框架自动处理搬运和同步）→ 硬件；高阶 API 通过调用基础 API 实现。
• Regbase：核函数 → UB 指针 → LoadAlign（手动搬运到寄存器）→ 计算 → StoreAlign（手动搬回 UB）；Reg API 直接调用编译器 BuiltIn API，高阶 API 和基础 API 也可以向下调用 Reg API。

MemBase vs RegBase

维度	MemBase（基础 API）	RegBase（Reg API）
数据载体	LocalTensor（UB）	RegTensor（VF Reg）
中间结果	必须回写 UB	可在寄存器中连续消费
单次处理粒度	任意长度（硬件内部 tiling）	一个 VL，需软件显式分块循环
数据搬运	DataCopy（MTE2/MTE3）	LoadAlign / StoreAlign（UB ↔ Reg）
Mask 控制	count 参数自动处理	显式 MaskReg 寄存器
寄存器复用	硬件决定	软件可控（dst=src 时直接复用）
性能上限	受 UB 往返与 MTE 流水限制	减少 UB 往返，提升指令并发
编程复杂度	较低	较高，需理解寄存器与调用层级

典型代码流程

一个标准的 REGBASE 算子 Kernel 通常仍采用三段式结构（CopyIn / Compute / CopyOut），只是 Compute 段被细化为"Load → 计算 → Store"：

// 1. CopyIn: GM → UB（与 MemBase 相同，用 DataCopy / TPipe 队列）
// 2. Compute: UB → Reg → 计算 → Reg → UB
template <typename T>
__simd_vf__ inline void AddVF(__ubuf__ T* dstAddr, __ubuf__ const T* src0, __ubuf__ const T* src1)
{
    // 本函数在矢量运算单元上执行，内部使用 Reg API
    AscendC::Reg::RegTensor<T> vDst, vSrc0, vSrc1;
    AscendC::Reg::LoadAlign(vSrc0, src0);        // UB → Reg
    AscendC::Reg::LoadAlign(vSrc1, src1);
    AscendC::Reg::Add(vDst, vSrc0, vSrc1);       // Reg 内计算，中间结果不落 UB
    AscendC::Reg::StoreAlign(dstAddr, vDst);     // Reg → UB
}
// 3. CopyOut: UB → GM

完整调用链是：核函数（__global__ __aicore__）从标量控制流发起，通过 asc_vf_call 调用 __simd_vf__ 函数，进入 VF Reg 层做寄存器级计算。

典型应用场景

多步融合计算是 REGBASE 收益最大的场景。 当以下条件同时满足时，建议考虑 REGBASE：

• 目标设备支持 Reg 编程（Ascend 950PR / Ascend 950DT）；
• 核心耗时集中在连续矢量计算上（如 Cast → Mul → Add → Cast 这种链式结构，中间结果可全部留在寄存器）；
• 中间结果在 UB 中存在明显反复读写；
• 追求极致性能，可接受更高代码复杂度。
  此外还包括：计算密集型算子（Exp / Ln / Sqrt / Div 等指令周期长的运算）、双宽场景（int64 运算需要 RegTraitNumTwo 把 2 个 VL 拼成 2×VL 的逻辑寄存器）、非连续访问（Gather / Scatter、Block Strided Load 等 MemBase 难以高效表达的模式）。
  对应的，单步简单算子（纯 Add/Sub）、数据搬运本身是瓶颈（mte2_ratio > 90%）的算子、团队不熟悉寄存器模型的场景，用 MemBase 更合适；而且在开放寄存器编程的硬件上，Memory 矢量接口内部也会做软件封装的 Reg 矢量计算，以保持与前序代际硬件的兼容。

小结

REGBASE 与矢量单元、矢量寄存器是"编程模型 — 硬件执行单元 — 硬件存储资源"三位一体的关系：矢量单元（AIV）是执行载体，矢量寄存器（VF Reg）是新增的一级私有存储，REGBASE 则是让开发者直接驱动这两者的编程接口。它不是对 MemBase 的替代，而是给追求极致性能的开发者开放的一条更贴近硬件的路径——代价是更高的使用门槛，以及显式承担同步、分块和寄存器分配的管理责任。

一、950PR 新增的编程能力：SIMD/SIMT 混合编程模型

昇腾 950PR（以及 950DT）在架构层面引入了SIMD/SIMT 新同构设计，这是本次代际最核心的编程模型变化：

• 此前（910B/910C 及更早）：AI Core 仅支持 SIMD（单指令多数据）执行模型，编程基于 Ascend C，对 CUDA 生态不友好。
• 950PR 开始：向量计算单元同时支持 SIMD 和 SIMT（单指令多线程）两种并行模型，可以处理线程块、线程束、内核启动等类 CUDA 原生功能。
• 配套软件栈 CANN Next：新增了 CUDA 兼容的编程抽象，允许开发者把现有 CUDA 代码以较低成本迁移到昇腾平台。
  混合编程的定位是以 SIMD 为主（承担 90% 以上的密集算力）、SIMT 为辅（专门应对复杂控制流、离散访存等不规则场景），这也是字节跳动、阿里愿意下单的软件侧关键原因。

二、950PR 上跑 RAG：全管线 NPU 化

虽然 RAG 不是"编程方式"，但昇腾 CANN 生态确实围绕 RAG 场景做了完整的适配，950PR 是首个原生为这类场景优化的推理芯片：

RAG 环节	昇腾上的支持情况
Embedding 模型（如 bge-m3、bge-large-zh）	NPU 推理，ATB 加速库支持
向量检索	可将向量库存于 NPU 显存，用 Cube Unit 做余弦距离 MatMul，避免 CPU-GPU 数据搬运
Reranker 精排	bge-reranker 等模型在 NPU 上运行
LLM 生成	Qwen3、DeepSeek V4 等已全面适配
编排框架	LangChain、LlamaIndex 等通过 torch_npu 无缝集成
950PR 针对 RAG 这类场景的硬件优化点包括：128 字节 Sector-Cache（访存颗粒度从 512B 降到 128B，小算子访存效率提升 4 倍，对离散的向量检索访问特别有利）、112GB 大容量 HBM（可在片内放下更大的上下文和向量库），以及PD 分离架构（Prefill 阶段是 950PR 的主打场景）。有实测数据显示，同一段 128K 上下文的 RAG 检索，在 950PR 上延迟可稳定在 38ms 左右。

三、小结

"950PR 新增 RAG 编程方式"这个说法不太准确——更准确的表述是：

1. 编程模型层面：950PR 新增了 SIMD/SIMT 混合编程，并通过 CANN Next 提供了 CUDA 兼容能力，这是吸引开发者的核心变化。
2. 应用场景层面：昇腾 CANN + ATB + MindSpeed 的软件栈完整支持 RAG 全管线在 NPU 上运行，950PR 的硬件特性（大 HBM、细粒度访存、PD 分离）恰好对 RAG 推理的 Prefill 和检索环节做了针对性优化。
   如果你是想在 950PR/Atlas 350 上部署 RAG 应用，可以直接参考昇腾社区提供的 RAG 推理管线示例（基于 CANN + bge 系列 + Milvus/FAISS + Qwen/DeepSeek 的组合），不用从算子层写起。