自研加速器芯片背景知识补充

设计验证

设计验证（Design Verification，DV）指的是：在芯片流片前，用仿真、形式验证、emulation 等方法确认 RTL 设计行为符合规格。具体方法包括：RTL simulation、UVM testbench、CModel/golden reference 对比、assertion/formal verification、coverage、emulation、FPGA prototype、log/waveform/debug flow 等等。

软件仿真

软件仿真模型，指的是用完整的软件栈，把软硬件系统串起来，让同一套 workload 可以在软件模型、RTL/Emulation 或真实芯片等不同后端上都能运行；其作用一方面是做架构探索，另一方面是在没有硬件的情况下对架构和软件进行验证。以下是一个以 C Model 为模型实现方式的软件栈架构框图示例。

其中各层的职责分别如下：

• Framework 层包含 PyTorch、Triton 等框架入口，以及 Model Layer、DL Model 级别的测试。
• Arch Tests/Library 层包含 Arch/Micro Arch 级别的功能和性能测试，以及算子库的测试。
• Compiler 负责编译 Host 侧和 Device 侧代码，生成 executable 或 fatbin。
• Driver/runtime 向上提供驱动（Kernel Mode Driver, KMD / driver API）和运行时（User Mode Driver API, UMD / Runtime API），向下连接用户态驱动、内核态驱动、仿真驱动和真实硬件。
• CModel 对 Driver、QEMU、Trace Player 提供统一的 model interface，Arch Model 和 ESL Model 都挂在这个接口后面。
• Emulator 和 Chip 是更靠近真实硬件的后端，前者用于硬件仿真或 emulation，后者对应真实芯片。

软件仿真模型分为 Arch Model、Function Model 和 Performance Model 三类，在上述软件栈框图中，将 Function Model 和 Performance Model 合并为 ESL Model。ESL Model 是一种高抽象度的功能模型，用于在芯片设计的早期验证系统架构，而不描述具体电路，比如 Xilinx 的 HLS 就是 ESL 设计方法中的一种实现方式。

• Arch Model 主要模拟编程接口和 ISA 计算执行，需要保证计算逻辑、数值精度和硬件一致，但会尽量省略硬件模块内部细节。它开发周期短，调试方便，适合做架构级功能验证和细粒度模型的 golden reference。
• ESL Model 更接近硬件模块级结构，比 Arch Model 更精细。Function Model 关注硬件模块级行为和全系统功能仿真，目标是做硬件验证的 golden reference，也给底层软件开发提供一个硬件模拟层。Performance Model 关注和性能相关的硬件模块级时序行为，性能无关或影响较小的细节可以抽象掉，用来加快仿真速度。

RTL 仿真

RTL 仿真分 Simulator 和 Emulator 两种，两种仿真环境都可以进行波形抓取来进行细粒度的问题定位和性能分析，区别在于前者是用软件解释或编译 RTL 来执行，后者是将 RTL 映射到专用硬件中执行。Simulator 的主流运行平台包括 Synopsys VCS、Cadence Xcelium 和 Siemens Questa；Emulator 的主流运行平台包括 Cadence Palladium、Synopsys ZeBu 和上海合见的 UVHS。Simulator 更灵活更方便，可以查看任意多的波形；Emulator 更快，比 Simulator 快出几个数量级，但价格也比 Simulator 贵出几个数量级。另外 Emulator 本质上也是在 FPGA 上运行，但相较于做 Prototype 的通用 FPGA，调试能力更强，更适合做大规模 SoC 的系统级验证工作。

除了仿真模型本身之外，还需要配套的测试集。测试集是一套可编译运行、能判断 pass/fail 的验证 workload。同一批测试需要能够运行在 Arch Model、Function Model、Performance Model、RTL Simulation、Emulation 和 Chip 这些不同后端上，这样才能对齐 workload 来验证不同的后端，测试集自身的正确性则通过 CPU/GPU reference 或 SelfCheck 交叉验证。相比 UVM 随机激励，测试集更接近真实软件和真实 workload。

协同验证

这里的协同验证，指的是将软件仿真与 RTL 仿真在同一批 case 下进行结果对齐的过程，是设计验证的一种。

协同验证的工作流如图，重点是验证某个模块或子系统的接口行为是否和 ESL Model 一致。其中 Sequence 产生测试序列，Driver 把同一份激励送到 DUT、SV Model 和 ESL Model。Monitor 采集输出，Scoreboard 做对比。这里的 ESL Model 可以作为 golden reference，也可以输出 interface dump 给 RTL testbench 当激励。

协同验证的方法有 Co-Sim 和 Vector-Sim 两种。Co-Sim 是 DV 侧的 testbench / UVM sequence / 验证环境来生成激励，同时给 RTL 和 CModel；而 Vector-Sim 是软件模型先把 case 用 Vector Dump 的方式打包，DV 再拿 dump 的包去驱动 RTL。

Vector Dump 指的是软件仿真跑过一个 case 之后，把用于在 RTL 环境下复现的输入激励 + 期望结果 + 元数据落盘成 RTL testbench 可读取的软件包的过程，其作用是把软件模型的中间状态或接口数据导出，用来和 RTL 仿真的结果对齐，方便定位到模块接口、指令执行或数据搬运的具体差异。

host-device 异构计算执行模型

Heterogeneous Computing（异构计算）指 CPU、GPU、NPU、FPGA 等不同处理单元协同工作。host-device 模型是其中最常见的一种，现代 GPU 以及各类自研 AI 芯片基本都是典型的 host-device 异构计算。CPU 作为 host 侧，负责组织、调度、加载、提交和同步；GPU 或 AI 芯片作为 device 侧，负责高吞吐张量计算；runtime、driver 和 firmware 把两边接起来。

执行模型本质上是解决软件如何驱动硬件计算的一套抽象，涵盖程序如何变成目标设备上实际运行的指令的全过程，是 ISA + ABI + OS 调度模型的集合。所谓可编程的功能，也需要通过具体的执行模型来实现。一段高级语言编写的代码要运行在异构计算设备中，执行模型大致如下：

模型代码: PyTorch / vLLM / TensorFlow / JAX / ONNX Runtime
-> 框架把模型拆成 OP / graph
-> 后端选择每个 OP 的执行设备: CPU、CUDA GPU、TPU、NPU 或自研 AI 芯片
-> dispatcher 或编译器决定每个 OP 怎么执行
-> runtime 分配显存、管理 stream/event、提交 kernel
-> driver 负责命令队列、DMA、中断和权限，把命令发送到设备
-> firmware / command processor 调度设备侧任务
-> 矩阵单元、向量单元、DMA、SRAM、HBM/DRAM 完成实际计算
结果回到框架 tensor

CUDA 编程模型是典型的 host-device 异构模型，包含主机端代码和设备端代码。核函数通过 grid 和 block 组织线程。一个数组相加例子如下：

#include <cstdint>
#include <cuda_runtime.h>

__global__ void add(const int* A, const int* B, int* C, int n) {
    uint32_t id = blockDim.x * blockIdx.x + threadIdx.x;
    uint32_t step = blockDim.x * gridDim.x;

    for (uint32_t i = id; i < n; i += step) {
        C[i] = A[i] + B[i];
    }
}

#define N 1024

int main(int argc, char* argv[]) {
    size_t dsize = N * sizeof(int);
    int host_A[N];
    int host_B[N];
    int host_C[N];

    int* data_a = nullptr;
    int* data_b = nullptr;
    int* data_c = nullptr;
    cudaMalloc((void**)&data_a, dsize);
    cudaMalloc((void**)&data_b, dsize);
    cudaMalloc((void**)&data_c, dsize);

    cudaMemcpy(data_a, host_A, dsize, cudaMemcpyHostToDevice);
    cudaMemcpy(data_b, host_B, dsize, cudaMemcpyHostToDevice);

    dim3 grid{16};
    dim3 block{2};
    add<<<grid, block>>>(data_a, data_b, data_c, N);

    cudaMemcpy(host_C, data_c, dsize, cudaMemcpyDeviceToHost);

    cudaFree(data_a);
    cudaFree(data_b);
    cudaFree(data_c);

    return 0;
}

框架、算子和 kernel

框架接入是用户可见的入口，通过写 torch.tensor(..., device="my_accel")、model.to("my_accel")，或通过推理引擎，可以在目标设备上运行模型。Dispatcher 是框架内部的分派机制，用来根据 OP、dtype、device、layout 等信息选择具体后端实现。

OP（operator，算子）是框架语义层的“操作”。用户写的模型会被框架表示成一组操作，例如 matmul、softmax、reshape、attention。graph 是这些 OP 以及数据依赖关系组成的计算图。

kernel 是设备上真正执行的一段程序，由程序员声明并给定具体 shape、dtype、layout、tile 切分和并行策略等参数，在设备上由多个线程并行执行。在不同的目标设备上，kernel 的线程模型、内存层次和指令集可能完全不同。一个 OP 不一定对应一个 kernel。编译器可能把多个 OP 融合成一个 kernel（融合算子），一个 OP 也可能拆成多个 kernel。

OP 和 kernel 中间隔着图优化、layout 选择、tiling、memory planning、kernel selection、autotune 和 runtime 调度等，很多性能问题就藏在这个中间层：尽管语义上只是一个简单 OP，但落到设备上可能产生额外 layout conversion、低利用率 tile、同步等待或 HBM 往返等等。

典型算子可以粗略分几类：

类型	常见 OP	主要压力
矩阵计算	matmul、GEMM、BMM	矩阵单元吞吐、tile 利用率、HBM 带宽
逐元素计算	add、mul、activation、quant、dequant	向量单元吞吐、访存开销
归一化	softmax、RMSNorm、LayerNorm	reduce、特殊函数、数值稳定性、片上缓存
形状和布局	reshape、view、transpose、contiguous	layout 表达、是否产生真实拷贝
注意力	attention、RoPE、KV cache 访问	矩阵计算、softmax、片上存储、显存带宽
稀疏/路由	embedding、MoE、expert parallel	随机访存、通信、负载均衡
通信	allreduce、allgather、reduce-scatter、send/recv	多卡互联带宽、拓扑、同步

编译器相关

编译器完成源代码到目标平台二进制文件的转化，是最难的工作之一。编译器前端要接 PyTorch、JAX、TensorFlow 或 ONNX 等框架，中端可能使用 StableHLO、MLIR、XLA 或自研 IR，后端再做 fusion（算子融合）、constant folding、shape specialization、layout optimization、memory planning、tiling、lowering、寄存器分配、指令调度、代码生成和 autotune 等。

编译器会基于底层硬件指令（ISA 操作语义）封装一系列 builtin 函数。程序员编写源代码并调用 builtin 函数来驱动硬件完成计算。Builtin 是编译器或语言内置函数，通常对应特殊硬件操作或低层能力。

ISA（Instruction Set Architecture，指令集架构）定义处理器能听懂哪些指令，规定了有哪些指令，有多少寄存器，每条指令的格式，地址空间，原子操作，内存模型等等，可能围绕矩阵乘、向量计算、load/store、DMA、同步、低精度转换等进行设计。对 AI 芯片来说，ISA 不一定像通用 CPU 那样暴露给普通程序员。它可能更多服务于编译器、手写 kernel 工程师和算子库。

算子开发依赖两类基础能力：runtime 提供的基础功能 API，以及 compiler 提供的 Tile Core 和 RV Core API。当前 kernel 开发主要使用 Tile Core builtin 指令集合和 RV Core builtin 指令集合。为了提高执行效率，有时会手写汇编，也就是直接编排 Tile Core 和 RV Core 的指令。手写汇编的目的通常有两个：一是调整指令顺序，让不同功能单元持续有活干；二是展开寄存器使用，减少数据依赖。

类似于 __global__、__device__、__shared__ 的称为编译器标识符，下面是 CUDA 中常见的三个编译器标识符：

• __global__：声明一个 CUDA kernel 函数，由 host 端用 «<…> » 启动，在 GPU device 上执行。
• __device__：声明一个 device 端函数或变量，只能在 GPU 端代码中调用或访问。
• __shared__：声明 线程块内共享内存变量，同一个 block 里的线程都能访问，生命周期限于该 block。

disable_tile_scheduling 可以作用在 global 或 device 函数上，编译器会尽量让这些函数里的 Tile 指令按源代码顺序出现在汇编中，这通常叫 Tile 指令保序。

builtin 指令是编译器内建指令，属于源代码层面的特殊接口。汇编指令则是 ISA 层面的具体指令。builtin 的价值在于：让源代码能表达普通 C/C++ 语法无法表达、但硬件确实支持的能力。

可编程相关

Runtime 是高层运行时库，如 CUDA Runtime API，面向程序员提供了内存分配 malloc、设备之间的内存拷贝 memcpy、kernel 启动、module load、stream 任务队列、event marker 等功能，面向硬件通过调用 Driver 提供的接口实现对硬件的操作。现实路线通常是先做 CUDA-like 的窄 API，保证能分配内存、拷贝数据、加载 ELF、launch kernel 和同步，然后再补 stream、event、graph capture、异步执行和内存池等功能。

Driver 是直接管理硬件的软件，功能包括设备初始化、内存映射、命令队列提交、DMA、中断、上下文切换、binary 装载等。通过 Driver 的接口把指令给到硬件的 control 部分，硬件根据指令中各个位置的信息解码出 ISA 语义上的各个参数（shape、地址、模式等），并 trigger 硬件执行。

Memory Model（内存模型）是针对编程语言的概念，说明了有哪些 memory，并定义不同内存空间、cache、shared memory、SRAM、HBM 和主机内存之间的访问规则、生命周期、同步规则、搬运规则。内存模型描述了编程语言如何抽象和访问内存；存储器硬件只提供线性的地址空间，不同编程语言需要在这个线性空间上建立自己的软件抽象以便于程序员使用。C 抽象为字节序列+指针、Java 抽象为对象+堆、Python 抽象为对象+垃圾回收堆。内存模型的设计初衷是屏蔽硬件差异，保证程序在不同平台之间的可移植性。

C/C++的内存模型如下，在编译生成的 map 文件中可以找到对应的概念。

栈、堆、全局区（data/bss）本质上都是地址空间中的一段区域，最后都会映射到真实内存里，区别主要是管理方式。局部变量通常在栈上，函数返回后自动回收；malloc/new 申请的对象在堆上，需要手动或由运行时管理释放；全局变量和静态变量通常在 data/bss。之所以会产生栈访问，常见原因是局部变量无法全部放进寄存器，发生寄存器溢出后，编译器会生成保存到栈上再取回的指令。

存储硬件相关

L2B 是 L2 Buffer，由 SRAM 做成，挂在 PE 旁边作为本地存储区。它的作用是降低全局带宽需求、降低访问延迟，也减少功耗。L2B 更接近显式管理的数据缓冲区，可能由 DMA、编译器或调度器控制，不像 CPU cache 那样完全由硬件自动管理。加速器的数据访问模式通常更规则，很多访问可以提前知道，所以显式 buffer 往往比自动 cache 更可控。

L2C 是 L2 Cache，属于硬件自动管理的缓存机制。程序员按地址访问，不需要直接控制 cache 的交互。cache line 是 cache 读写、命中和替换的基本单位；不同硬件单元同时访问同一个 cache line 时，可能触发额外同步、失效、刷新和重载，性能会下降。

L2B 和 L2C 都在存储层级的第二级，所以都叫 L2。按离计算单元从近到远，可以粗略分成寄存器、L1、L2、L3 或 Global Buffer、DRAM/HBM。

interleave 是交错编址：把多个 DRAM 组合成一个更大的逻辑内存空间。这里的 interleave 粒度是 4KB，意思是每 4KB 轮流放到不同 DRAM 里。它的目的不是单纯扩大容量，而是提高带宽，避免单个 DRAM 变成瓶颈，并让访问更均匀地打到多个 DRAM 上。

地址对齐可以理解为硬件希望数据从固定大小的边界开始存放，而不是从任意地址开始。首地址对齐指数据块第一个字节的地址满足对齐要求。比如要求 32B 对齐，就表示首地址必须是 32 的整数倍。硬件访问数据通常一次读一整块；不对齐时，一次访问可能跨块，最后变成两次访问。Linear 搬运一次搬一个 RV Vector，要求 32B 对齐；Tiled 搬运一次搬整个 Tile 的一部分，要求 256B 对齐。