引言
模拟器技术的思想可以追溯到图灵机，艾伦·图灵在他 1937 年发布的专著中提出了图灵机的概念，对模拟技术进行了相关定义，这也是计算机模拟技术的理论基础。
因为模拟器可以帮助研发人员深入了解目标机的软件运行情况，包 括检测指令执行周期，内存使用，外设串口的正确性等。这些信息对硬件架构师，指令设计者，系统软件开发者来说有着重要意义。
一、模拟器需求分析
在传统的芯片研发流程中，设计某种芯片的周期往往很长。某一些模块或者子系统的功能，至少要等到FPGA验证阶段才可以对其进行分析，假如性能不合适，还需要推翻原来的架构重新设计，这毫无疑问增加了前期芯片开发的成本，拖慢了研发流程进度。
本文所研发的模拟器是一款AI芯片研发流程中的一个重要组成部分，主要是为了解决以下三个研发流程难点，提高芯片总体的研发效率，缩短研发流程周期而策划研发的。
一是验证AI指令集的正确性，识别并执行不同的指令类型，验证指令流以及芯片设计的功能；二是支持内部指令流执行信息查阅，从而加速相关芯片架构的探索和RTL设计的快速迭代，提升设计效率；三是适配验证多种并行算子，保证AI算法在芯片运行正确，使其能够满足AI处理器研发的各种需求。
该模拟器的开发意图是为了方便开发团队可以避免硬件设计流程的繁琐，方便前期架构的探索，用软件的方法，采用“串行”执行的软件语言去描述“并行”执行的芯片中各个模块的行为，并根据结果对硬件架构和工具链研发进行调整，因此本模拟器的研发重点是实现芯片功能上的模拟，并没有提供精确到时钟的仿真精度。
该模拟器的目标用户主要由处理器设计团队，工具链研发团队以及算法团队组成，他们对AI芯片的结构组成，SIMT的编程模型，并行计算的底层原理都非常了解，具备着相当专业的研发知识。
为了缩短研发周期，尽快的加速研发流程，本模拟器将采用解释型模拟策略，对处理器建立系统级功能模型。
对于此模拟器，虽然没有面向普通用户的通用交互界面和操作环境，但各个团队的研发人员都可以相当熟练地使用这款模拟器。
参与研发的该芯片有着相当独特的结构，强大的运算功能，大量的自研IP及最适配的软件优化。在研发过程中，因为时间性能等因素的影响，设计结构化繁为简，对指令执行模块以及指令模拟函数单元模块进行设计，实现了最为重要的AI指令集的模拟。
由于该模拟器要实现的指令编程基于SIMT(单指令多线程)思想，此架构会将线程的组织分为三个层次（Kernel-Block-Thread）。
Kernel对应处理器执行的基本任务，每个Kernel划分成多个彼此独立的线程组(Block)，每个线程组内部也包含有多个彼此独立的并行线程，他们之间可以实现细粒度的协同工作。
在这一过程中，模拟器需要提供模拟器初始化，线程维度数量控制，寄存器读写，初始地址等必要的接口控制，以便研发人员可以根据项目要求，随时对线程维度等常需变动的配置参数调用相关接口进行更改。
由于该模拟器的定位是芯片研发集成环境的一部分，不光需要对硬件设计团队的处理器指令功能进行对比，还需要执行将并行算子通过工具转化后的指令。
在这些过程中，模拟器还需要做到针对任意一特定线程的内部指令流执行信息查阅，从而针对性的验证指令执行的正确性，编译器编译算子表达的正确性以及算子的功能正确性。因此，该模拟器在研发流程中主要做到如图1.1所示功能：
（1）提供方便的外部配置接口以便研发人员调用。
（2）正确的模拟SIMT架构，模拟实现多线程结构的地址和内存设计，在功能上实现大密度并行指令流执行操作。
（3）可以完整的模拟AI指令集中的所有指令，并验证每个指令执行的正确性，从而提供正确的仿真模拟结果。
（4）可以针对性的查阅任意线程的内部指令流执行信息，以便及时的进行相关方面的调试。
图1.1模拟器需求
以上就是模拟器的具体功能分解。对于本模拟来说，最重要的是对输入输出模块，指令执行模块，以及线程层级结构与存储层级结构的细节设计。
二、指令集体系结构
2.1指令集整体描述
本模拟器所采用的指令集是一种很典型的精简指令集（RISC），每条指令的位宽均为64比特。为了更方便的进行计算机的流水线操作，本指令集主要将指令分为ALU/SFU/MEM/TC/CTRL等几类指令编码格式，各类指令的编码格式主要如图1.2所示：
图1.2各类指令编码格式
图中缩略词含义如下：
@p：predicate，分支跳转控制位，判断分支条件发生时跳转与否。
opcode/op：指令的机器码。
ctrl1/ctrl2:指令控制位，用来解析确定指令的具体格式。
dst:目的寄存器索引
src:源寄存器索引。
rs1/rs2/rs3:操作数所在寄存器的索引。
rsvd:此域值保留，不对本次指令行为造成影响。
sf1/sf2:MEM类型指令格式里的地址偏移量。
imm:采用立即数作为源操作数。
模拟器目前所支持的数据类型主要包括：U32、S32、U16、S16、U8、S8、FP32和FP16这几类。
除此之外，一部分指令还可以支持将数据pack到一起的packed类型的数据，比如将两个fp16pack到一起，形成fp16x2等。
加减，乘加，绝对值减等等指令都会支持fp16x2的计算。而其他ALU和SFU指令的操作数为packed数据类型时，还可以通过ctrl2的某几位指定操作数的起始位置位于源寄存器值的哪个byte。
2.2指令分类
本模拟器共需模拟62条指令，其中数据处理指令35条，格式转换指令1条，矩阵运算指令3条，数据混洗指令1条，load/store指令15条，分支及系统指令7条。不同功能的指令，格式也会不同。按照类型来分，主要有以下几类：
（1）ALU指令
ALU指令指的是算数计算与逻辑运算类指令，包括加减乘除，比较，极限值，与或非，数制转换以及移位等指令。
根据格式要求，从源寄存器中取出两个操作数或三个操作数后，统一数制并进行计算。ALU指令编码格式主要由opcode、ctrl与操作数组成。
opcode代表具体的计算类型，ctrl用于选择数制、立即数、状态位等，ctrl2用于指定操作数偏移。
在ALU类指令中，格式转换指令比较特殊。它可以将一个操作数由src数制转换成dst数制，实现整型到浮点以及浮点到整型的转换。
转换过程可以指定saturation和rounding mode，假如dst数制为浮点型，还可以对输出结果做对齐。
除此之外，convert指令还可以将两个操作数做数制转换，并对结果做拼接。
（2）Special指令
Special类指令大致可分为以下三大类：SFU（超越）函数指令，TC(矩阵运算)指令，shuffle（寄存器值混洗）分配指令。这三类指令都分别拥有各自的编码格式，根据格式要求，从源寄存器中取出操作数后，统一数制并进行计算。
SFU指令编码格式与ALU指令相同，主要由opcode、ctrl与操作数组成。opcode代表具体的计算类型，ctrl位控制数制和状态位等，指令根据计算类型，将操作数进行计算后存入相应目的寄存器。
此外，SFU指令只会运行在thread级计算单元上。当指令的目的寄存器为thread寄存器时，源操作数可以来自thread寄存器，warp寄存器以及常数寄存器。
TC（Tensor core）指令主要是用来实现矩阵运算，它会根据指令规则，分别将两个操作数所在warp线程组内的所有线程的对应寄存器值取出，组成两个矩阵进行计算，并将结果输出相同warp下各线程的目的寄存器中。
如果源操作数是半精度浮点数，目的操作数为单精度浮点数，输出需要两个寄存器，那么溢出的输出将会顺延到加一索引的寄存器上面。
SHUFFLE指令格式与TC指令基本相同，它所实现的主要功能是将该指令所在warp下的所有线程中源寄存器的值按某种规则重新分配。
（3）MEM指令
MEM指令编码格式主要由opcode、atom、ctrl与操作数组成。opcode用于区分MEM指令类型，Atom有4比特位宽，用来区别此条指令是普通存储指令还是原子操作指令。
LD指令可以根据计算地址，将数据从存储器中读取至目的寄存器中，ST指令则可以将源寄存器值，根据计算地址保存至存储器中。LD/ST地址计算方法如下:
ܾ
（1-1）
(1-2 )
由（1-1）和（1-2）式可以得到指令的基地址base和偏移地址offset，结合译码出的立即数，可以推得目的地址如下：
(1-3)
݉݉ݏMEM指令可以同时运行在thread级计算单元和warp级计算单元上。对于LD指令来说，目的寄存器值只可以放入thread寄存器和warp寄存器；ST指令的源寄存器值可以来自thread寄存器，warp寄存器和常数寄存器。rs1操作数可以来自warp寄存器和常数寄存器，rs2操作数来自thread寄存器和常数寄存器。
由于处理器的设计基于SIMT思想，会将一个大的计算任务分解成为多个小的线程来计算，每个线程都拥有独立的索引和寄存器空间，因此MEM的目的地址还支持以下几种应用场景：
1.普通地址（warp/thread）+固定偏移量（imm）
2.普通地址（thread）
3.普通地址（warp/const）
4.普通地址（thread）+tid(threadid)
5.普通地址（warp）+tid(threadid)
6.Globalbase（warp）+tid（threadid）
MEM类指令除了LD/ST指令外，还有一类锁线程的原子操作指令。原子操作指令由连续的一条或多条atomicST和一条atomicLD指令组成。
atomicST用于设置操作数，atomicLD用于根据atom字段，以及前面设置的操作数，执行相应的操作。
（4）CTRL指令
CTRL指令指的是控制类指令，包括exit（中止），nop（气泡），loop（循环），jmp（跳转），branch（分支）等指令，主要起到逻辑控制的作用，在程序运行中非常重要，下面描述几个典型ctrl指令。
LOOP指令循环指令，主要由opcode、ctrl与操作数组成。ctrl低七位表示参与循环的指令数目，最高位用于选定立即数。rss，rsi，rse限制循环次数，值可来自warp寄存器，常数寄存器以及立即数。
JMP指令无条件跳转指令，支持绝对地址跳转和相对地址跳转，跳转值来自寄存器或立即数。其中Bit[0]:0代表用做地址偏移，1代表用做绝对地址；Bit[7]:0代表rs1作为地址，1代表imm作为地址。
BRANCH分支跳转指令，与无条件跳转指令类似，同样支持绝对地址跳转和相对地址跳转，跳转值来自寄存器或立即数。此外，branch会根据指令63位的@p以及ctrl位决定是否跳转。
EXIT指令终止目前的进程，标志当前程序的结束NOP指令“空指令”：仅当成一无实际意义的指令执行，并继续执行NOP后面的一条指令。用于清除flag以及配合其他指令起到一些辅助作用。
2.3线程组织及存储空间
由于处理器基于SIMT体系结构，因此相比起SIMD，它拥有独特的线程组织和寄存器空间。
如图1.3所示，线程的组织分为三个层次（Kernel-Block-Thread）。Kernel对应模拟器执行的基本任务。每个Kernel划分成多个彼此独立的线程组(Block)。
每个线程组内部包含多个线程，他们之间可以实现细粒度的协同工作。每个线程对应的指令代码相同，但根据参与运算的数据不同，输出结果也会相应不同。线程内部支持全部的C语法。
图1.3SIMT编程模型
Block和Thread均采用三维方式组织，每个thread，warp，kernel均有自己的存储空间。此外，对于每一个线程，可以通过特定的寄存器访问到自己对应的blockIdx.*，threadIdx.*值。
在模拟器需要实现的SIMT编程模型中，单线程内部执行的指令会做到严格保序；在多个线程执行指令时，指令之间执行顺序彼此独立，不会互相影响；当不同的线程访问同一片存储空间时，由于线程的优先级和仲裁规则未确定，此时的程序行为是未定义的。
由于不同线程访问同一片存储空间对于某些程序是必要的，此时需要引入线程间的协同机制。线程的协同需要编程人员调用intrinsic function来实现，模拟器也支持通过相关指令功能的实现。
每个线程能够访问的存储空间分为几个部分：registers、local memory、shared memory、external memory。其中寄存器和Local memory为thread独享空间，Shared Memory为Block内部共享，External Memory为整个kernel所共享。
除此之外，local memory，shared memory与external memory之间存在一定的地址映射关系。存储结构模型如图1.4所示：
图1.4SIMT存储空间模型
每个线程都有各自的寄存器空间regfile，目前主要包含线程寄存器thread regfile，线程组寄存器warp regfile，常数寄存器constant regfile这三类。Thread regfile为每个thread计算使用的通用寄存器，不同thread之间独立。
每个线程最大分配128个thread寄存器；Warp regfile和thread regfile在编程抽象上的地位相同，核心区别在于Warp级别的同步性；Const regfile包含两个部分，程序的形参和程序运行时的特殊参数。
图1.5寄存器编码格式
寄存器的编码格式如图1.5所示，其中低128位为thread寄存器；128-159位为warpregfile；160-191位为indirect warp regfile，其与128-159位warp寄存器一一对应，存储相应立即数；192-255位为constant regfile，是只读寄存器。
三、模拟器整体实现
模拟器的主要功能是模拟实现处理器功能行为。在深入研究学习了芯片的体系结构后，结合项目需求以及开发环境后，本设计最终采用结构清晰，仿真精度高，易于设计和调试的解释型模拟策略，使用system C语言，对AI芯片建立系统级功能模型。
图1.6模拟器框架图
本模拟器的整体框架如图1.6所示，主要由三大功能模块组成，分别是提供时钟和复位信号的gen模块、指令集核心模拟模块、存储模块。由于本模拟器受研发周期和需求限制，仅需对处理器建立功能模型，暂不精确至单个周期，gen模块仅是作为概念存在，待后续开发流程补充。
此外，模拟器一共需要五个文件，分别是二进制指令文件、配置文件、输入数据文件、输入数据地址文件、输出数据地址文件。
其中，二进制指令由程序经过编译器生成并存放进程序存储器中；配置文件作为内核信息参数表，含有一些模拟器配置参数，也会在模拟器初始化时存入存储器中；输入数据文件存放了将要放入存储器，参与模拟器运算的数据；输入输出地址文件中分别放有输入数据存放存储器地址和输出数据来源地址。
图1.7模拟器顶层工作流程图
模拟器的主要模块围绕着核心的指令执行模块工作，由于需要考虑到SIMT编程模型和AI指令的特性，对于线程级指令和线程组级指令，总体来说有如图1.7所示的流程：
（1）仿真平台时钟和复位信号的定义，复位、模拟器等功能模块以及存储模块的实例化。
（2）通过初始化函数，对程序计数器，寄存器，MEM等存储结构初始化。
（3）运行读取函数，将二进制指令文件，处理器配置文件及HBM输入数据并存至相应线程的寄存器及存储空间中。
（4）模拟器从存储模块读取配置文件，生成相应规模的维度结构，并以PC为单位，按照顺序遍历的方式在每一维中独立运行取值、译码、取值、执行、写回等指令执行模块。
（5）检测当前PC值是否为该线程可执行指令流末尾，若是，则等待其他线程结束后进行下一步；若不是，该线程将跳转至根据计算得到的PC处继续执行指令。
（6）运行输出函数，根据输出数据地址文件，将文件地址的数据输出至指定文件中。
（7）结束仿真。
在指令执行模块运行前，模拟器会将程序存储器中存放的二进制指令存入各线程独立的程序寄存器中。在取指阶段，各线程将根据PC，独立的从程序寄存器中读取二进制指令；在译码阶段，各线程通过译码函数分离二进制指令的
opcode位，ctrl位，状态位并将这些信息分别存放进各个寄存器中；在取值阶段，模拟器会根据译码获得的寄存器索引，将对应的寄存器值放入各线程临时变量中并送至执行模块；执行模块则会解析中间变量中的opcode，通过指令功能模拟函数实现对应的指令功能；在最后的写回阶段，各线程会将功能函数得出的计算结果根据索引存入相应寄存器中。
对模拟器要实现的SIMT指令集体系结构来说，源操作数和目的操作数可能来自于不同的寄存器层次。当各个线程指令的执行只涉及thread线程寄存器和常数寄存器时，由于各线程thread寄存器的独立性，各线程的指令执行过程互不影响。
Warp寄存器和thread寄存器在编程抽象上的地位虽然相同，但warp寄存器需要保存在线程组内的同步性。因此，与thread指令不同，对于warp寄存器参与的指令，模拟器需要在指令执行过程中对各个线程的warp寄存器进行相关同步调整。
接下来将分别对仅有thread寄存器参与的指令执行流程和有warp寄存器参与的指令执行流程进行分别说明。
（1）仅有thread寄存器参与的指令执行。
图1.8为仅有thread寄存器参与的指令执行流程，模拟器会依次遍历各线程组下的各个线程，并通过以下操作执行指令：
1、指令预取。根据PC，各线程从程序寄存器中读取二进制指令并存至中间寄存器中。
2、指令译码。确定运算类型，控制信息，源寄存器，目的寄存器，源地址，目的地址等信息。
3、指令取值。根据译码获得的寄存器索引，将各自线程中的thread寄存器值放入各线程临时变量中并送至执行模块。
4、指令执行。各线程通过指令功能模拟函数实现对应的指令功能。
5、数据写回。将指令执行得到的结果写入各线程thread寄存器中。
图1.8thread寄存器参与的指令执行流程
（2）有warp寄存器参与的指令执行。
图1.9为有warp寄存器参与的指令执行流程，模拟器会依次遍历各线程组下的各个线程，并通过以下操作执行指令：
1、指令预取。根据PC，各线程从程序寄存器中读取二进制指令并存至中间寄存器中。
2、指令译码。确定运算类型，控制信息，源寄存器，目的寄存器，源地址，目的地址等信息。
3、指令取值。根据译码获得的寄存器索引，将各自线程中的thread寄存器值或warp寄存器值放入各线程临时变量中并送至执行模块。
4、指令执行。各线程通过指令功能模拟函数实现对应的指令功能。
5、数据写回。根据配置参数在线程组内指定某一线程，将该线程指令执行得到的结果从中间寄存器写入该线程组内所有线程中相同目的索引的warp寄存器中。
不管是仅有thread寄存器参与的指令还是有warp寄存器参与的指令，模拟器都会以多重循环的方式，以顺序遍历的方式在线程组下的各个线程中执行指令；在线程组内部，模拟器也将按顺序遍历的方式，以从小到大的方式逐个执行指令。
Thread寄存器不存在数据冲突，因此对其工作流程不需调整；当warp寄存器作为源寄存器时，编译器会保证warp寄存器值的统一性，当warp寄存器作为目的寄存器时，模拟器可以在执行指令循环中通过配置指定参考线程来确保warp寄存器的同步性，提高算法程序适配性和性能。
图1.9warp寄存器参与的指令执行流程
总结
本文先从此次模拟器开发的使用场景、功能需求开始，确定了模拟器实现策略和开发方向；其次，对将要模拟的指令集的整体情况、编码格式、指令类型和线程、存储体系结构进行了描述；
最后通过以上信息，给出了模拟器的整体框架、层级结构和基本工作流程，基本说明了本模拟器的核心结构和设计思路。