引言
模拟器技術的思想可以追溯到圖靈機，艾倫·圖靈在他 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寄存器參與的指令執行流程
總結
本文先從此次模拟器開發的使用場景、功能需求開始，确定了模拟器實作政策和開發方向；其次，對将要模拟的指令集的整體情況、編碼格式、指令類型和線程、存儲體系結構進行了描述；
最後通過以上資訊，給出了模拟器的整體架構、層級結構和基本工作流程，基本說明了本模拟器的核心結構和設計思路。