1、前言&本文重點
在 GPGPU 顯得愈發重要的今天，僅憑 nVidia， AMD 提供的程式設計介面來了解 GPU 未免顯得太單薄了些。時至今日， GPU 內部如何執行一條指令的對程式設計師來說依然是透明的、不可見的。為了達到程式效率最最佳化的目的，就必須要對 GPU 工作過程有一定了解和認知，故本文以前人成果 （GPGPU-SIM）為例，試分析 GPU 執行指令的具體過程。
/*
本文最開始的目的是翻譯一下GPGPU-SIM的文件，中文文件在網路上並沒有看到很好的，有，還剩一些訪存的部分，這周爭取讀完。
*/
2、基礎知識
2。1 CUDA程式設計簡介
以下參考《CUDA C Programming Guide》
CUDA 是 NVIDIA 釋出的在其統一架構 GPU 上進行
通用程式設計
的並行程式設計環境。 CUDA 軟體環境包括一組 Runtime API、一組裝置驅動函式、以及一個庫檔案。它們的層次關係如圖下圖所示。
由上圖我們可以很清晰地看到，CUDA 驅動庫函式直接控制底層硬體結構Runtime 函式是對驅動函式的封裝。應用程式可以直接呼叫底層驅動函式，也可以透過呼叫 Runtime函式間接操作底層硬體。
CUDA 還包括有多個數學工具庫——諸如cuBLAS， cuFFT， cuRAND 。英偉達還提供一個被稱為 nvcc的編譯器。CUDA 所使用的程式語言基於 C/C++，並在 C/C++語言上進行了一系列的擴充套件，拓展主要包括以下四個方面：
用來表示函式是在主機 CPU 還是在裝置 GPU 上執行的關鍵字，
__global__
表示該函式為核心函式，只能在裝置上執行，
__device__
表示 在裝置上執行的非核心函式，
__host__
表示只能在主機 CPU 上執行的函 數；
用來表示變數位於 GPU 哪一種記憶體空間中，
__constant__
表示該變數位於常亮儲存中，
__shared__
表示該變數位於共享記憶體中；
指定核心函式的
並行度
，也就是 Grid 、 Block 的維度，例如
gridDim
，
blockDim
；
用於儲存 Grid 和 Block 的維度資訊和執行緒的索引標號，例如
blockIdx。x
，
blockIdx。y
，
blockIdx。z
，
threadIdx。x
，
threadIdx。y
，
threadIdx。z
。
2。2 CUDA 執行緒組織模型
2。2。1 抽象的執行緒組織
以下參考《PARALLEL THREAD EXECUTION ISA：2。3 Memory Hierarchy》
具體到硬體的執行部分，每個CTA所能包含的執行緒數是有限的，與此同時，每個核心函式均是被組織成多個CTA，這些個CTA是同時執行的，因此，一個核心函式可以啟動數量龐大的執行緒，不幸的是，不同CTA之間的執行緒無法通訊或同步，但相同CTA之間的執行緒可以同步或通訊。
每個執行緒可以訪問不同的資料空間，但略有限制：每個執行緒之內有自己的私有記憶體空間，每個執行緒塊之內有自己的共享記憶體，對塊內所有執行緒可見，與整個塊具有相同的生命週期，但所有執行緒均可以訪問全域性記憶體。
Tip： A cooperative thread array （CTA） is a set of concurrent threads that execute the same kernel program。 A grid is a set of CTAs that execute independently。
2。2。2 實際的執行緒組織
以下參考《CUDA C程式設計權威指南》
下面以二維矩陣元素為例，介紹 CUDA 軟體中執行緒組織形式
首先我們知道，在 CUDA 中矩陣是行優先儲存的，如下圖
前文介紹了，CUDA 中執行緒索引、塊索引為拓展字，因此我們可以直接拿來計算矩陣的元素位置
ix = threadIdx。x + blockIdx。x * blockDim。x
iy = threadIdx。y + blockIdx。y * blockDim。y
所以不難推知，當前執行緒塊內、當前執行緒操作的矩陣元素索引為
idx = iy * nx + ix
2。3 CUDA 機器模型
2。3。1 淺談三種並行模型：SIMD、SIMT、SMT
以下參考
《SIMD < SIMT < SMT： parallelism in NVIDIA GPUs》
《PARALLEL THREAD EXECUTION ISA：3。 PTX MACHICE MODEL》
首先釐清概念：
SIMD：單指令多資料，首先獲取多個數據，同時使用一條指令處理
SMT：同時多執行緒，不同執行緒之間的指令可以並行執行
SIMT：二者折中方案，單指令多執行緒，執行緒內部執行相同指令，但比SIMD更靈活，比SMT效率更高
其次，對比 SIMT 與 SIMD ，上文說到，SIMT 比 SIMT 更靈活，其主要體現在以下三點
1。 單指令，可以訪問多個暫存器組。（聯絡上文，每個執行緒有自己的暫存器）
單指令，多種定址方式。
單指令，多種執行路徑
（每組執行緒中，如果出現分支指令，則不同執行緒之間序列執行，直到分支指令執行完畢，每組執行緒繼續並行執行相同指令，下文會提供一種分支指令預測機制）
最後，對比 SIMT 與 SMT，上文說到，SIMT比SMT效率更高，主要體現在 SIMT 可同時執行的執行緒更多、暫存器更多這兩點：
足夠多的執行緒，可以獲得足夠高的吞吐率
一方面延遲是竭力避免的，另一方面暫存器的價格是可以接受的。
2。3。3 什麼是Warp
Warp是SM的基本執行單元。一個Warp包含32個並行thread，這32個thread執行於SIMT模式。也就是說所有Thread執行同一條指令，並且每個Thread會使用各自的資料執行該指令。
不難推知，每個block中Warp數量可以很簡單地推算出來，目前 nVidia 的GPU中 WarpSize = 32。
一個 Warp 中的執行緒必然在同一個 Block 中，如果 Block 所含執行緒數目不是 WarpSize 的整數倍，那麼多出的那些thread 所在的 Warp 中，會剩餘一些 inactive 的 thread，也就是說，即使湊不夠 Warp 整數倍的thread，硬體也會為 Warp 湊足，只不過那些 thread 是 inactive 狀態，需要注意的是，即使這部分thread是inactive的，也會消耗SM資源。
一個Warp內部產生的分支分歧問題，將在下文詳述。
2。4 CUDA 組合語言舉例分析
2。4。1 回顧CUDA 程式碼格式
以下參考《CUDA C Programming Guide》
函式的宣告，需要指出其執行的具體位置，是GPU還是CPU？
__global__ void foo（。。。）
// runs on GPU， callable from CPU
__device__ void bar（。。。）
// function callable from a GPU thread
需要制定 Grid ，block 的大小，以啟動核函式，用三個尖括號括起：
<<
foo<<<500， 128>>>（。。。）；
// 500 blocks， 128 threads 來啟動核函式
在核函式內部，需要透過計算來得到當前執行緒的ID
dim3 threadIdx； dim3 blockIdx； dim3 blockDim
2。4。2 CUDA 程式碼示例
以向量加運算為例，A［1。。N］ + B［1。。N］ = C［1。。N］
如果是用 C 語言書寫，只能用for迴圈實現
void vecADD_serial（
const int* a，
const int* b，
int *c，
const int n）
{
for （int i = 0； i < n； ++i）
c［i］ = a［i］ + b［i］；
}
但我們現在嘗試用 CUDA 語言書寫，每個元素均用一個執行緒操作，其核心函式如下所示：
__global__ void add_vectors（
const int* a，
const int* b，
int *c，
const int n）
{
const int idx = blockIdx。x * blockDim。x + threadIdx。x；
if （idx >= n） return；
c［idx］ = a［idx］ + b［idx］；
}
主函式大意如下所示，這裡需要申請GPU記憶體，交換資料CPU -> GPU ，啟動核心，傳回資料 GPU->GPU：
int main（） {
… // omitted： allocate and initialize memory
// Invoke parallel kernel with 256 threads/block
int nblocks = （n + 255） / 256；
add_vectors<<
… // omitted： transfer results from GPU to CPU
}
2。4。3 CUDA PTX分析
以下參考 《Demystifying PTX Code》
承接上文，透過
nvcc yourname。cu ——ptx
生成的組合語言如下所示：
//
// Generated by NVIDIA NVVM Compiler
// Compiler built on Sun May 18 04：44：51 2014 （1400399091）
// Driver 331。79
//
。version 3。0
。target sm_21， texmode_independent
。address_size 32
//檔案以nvcc編譯器的資訊註釋作為開頭，緊接著跟著三行：
//PTX ISA版本
//目標架構，計算能力
//使用的地址模式
。entry add_vectors（
。param 。u32 。ptr 。global 。align 4 add_vectors_param_0，
。param 。u32 。ptr 。global 。align 4 add_vectors_param_1，
。param 。u32 。ptr 。global 。align 4 add_vectors_param_2，
。param 。u32 add_vectors_param_3
）
//接下來是 。entry 指令指引的 kernel 函式入口
//下面是四個引數
//此Kernel函式的引數組指標，和一個32位整型變數（輸入和輸出向量的數是三個指向32位整型全域性記憶體中的長度）。
//每個引數以 。param 偽指令開頭，接著是其資料型別 。u32 （最佳化成了無符號整形）
//。ptr 引數為指標型別
//。global 資料都在全域性記憶體中
//。align 4資料對齊的方式，
//本例中是4位元組對齊
//注意，常量和
非常量指標
在這裡沒有區別
{
。reg 。pred %p<2>；
。reg 。s32 %r<21>；
//這裡是暫存器的定義
//以偽指令 。reg 開頭，暫存器名字是以 % 作為字首。
//有。pred指令的用來條件分配，比如分支指令。由於PTX是中間語言，因為暫存器的定義是虛擬的，不一定完全 和硬體暫存器是一對一的關係。
//一組包含N個虛擬暫存器的暫存器組可以用
ld。param。u32 %r9， ［add_vectors_param_3］；
//ld。param指令則將函式引數複製給了暫存器。因為絕大多數PTX指令都不能直接操作函式引數。所以第四個引數 傳遞的是地址，因此需要新增［ ］中括號來獲取其資料。
mov。u32 %r5， %envreg3；
mov。u32 %r6， %ntid。x；
mov。u32 %r7， %ctaid。x；
mov。u32 %r8， %tid。x；
//接下來，一些特殊的數值被複製給了GPU暫存器
//envreg3： 由驅動定義的、只讀特殊暫存器
//ntid。x：每個CTA的x維度的執行緒數量，相當於 get_local_sizeo（0）
//ctaid。x：grid裡的CTA識別符號，相當於 get_group_id（0）
//tid。x ：CTA x維度的執行緒號，相當於 get_local_id（0）
add。s32 %r10， %r8， %r5；
mad。lo。s32 %r4， %r7， %r6， %r10；
//首先是一個加法 add
//R10 = envreg3 + tid。x；
//其次一個乘法&加法 mad ：
//R4 = ntid。x * ctaid。x + tid。x = get_global_id（0）；
//PTX 文件中指出，envreg3 由驅動負責，可以不理會
setp。lt。s32 %p1， %r4， %r9；
@%p1 bra BB0_2；
//接下來就是本例中唯一的一個條件指令。
//setp指令是指比較 r4 （當前執行緒 ID ）是否比 r9（陣列長度）小（lower than，。lt指令）
//以此設定謂詞p1。
//@指令則判斷p1，如果p1是true，執行分支BB0_2
//bra指令，注意bra指令的目標一定要是label或者指向label的暫存器
//如果p1是false，則直接執行後面的程式碼。
//本例中，分支指令後僅有
ret指令
表示當前分支的結束。
ret；
BB0_2：
//實際的運算部分指令位於標籤 BB0_2 之下，總體包括
//計算資料指標
//從全域性記憶體中讀取資料
//實際的運算
//儲存資料到全域性記憶體中
shl。b32 %r11， %r4， 2；
//r4 暫存器內容，左移兩位（*4）存入r11，左移不考慮符號位
ld。param。u32 %r18， ［add_vectors_param_0］；
//讀取第一個函式引數，向量A起始地址，ld。param 讀引數
add。s32 %r12， %r18， %r11；
//相加得到當前執行緒 ID 操作的資料地址，存入%r12
ld。param。u32 %r19， ［add_vectors_param_1］；
add。s32 %r13， %r19， %r11；
//上同，獲得向量B的當前資料位置，存入%r13
ld。global。u32 %r14， ［%r13］；
//讀取向量B當前欲運算元，地址在r13中，ld。global 讀全域性記憶體
ld。global。u32 %r15， ［%r12］；
//讀取向量B當前欲運算元，地址在r12中
add。s32 %r16， %r14， %r15；
//實際的相加指令
ld。param。u32 %r20， ［add_vectors_param_2］；
add。s32 %r17， %r20， %r11；
st。global。u32 ［%r17］， %r16；
//上面三條指令，得到向量C的當前儲存地址，存數
ret；
}
2。5
記分牌
演算法扼要
2。5。1 簡介
記分牌是一集中控制部件，其功能是控制資料暫存器與處理部件之間的資料傳送。在記分牌中儲存有與各個處理部件相聯絡的暫存器中的資料裝載情況。當一個處理部件所要求的資料都已就緒（裝載完畢），記分牌允許處理部件開始執行。當執行完成後，處理部件通知記分牌釋放相關資源。所以在記分牌中記錄了
資料暫存器
和多個處理部件狀態的變化情況，透過它來檢測和消除或減少資料相關性，加快程式執行速度。
2。5。2 方法
儘可能提早指令的執行。當一條指令暫停執行時，如果其他後繼指令與暫停指令及已發射的指令無任何相關，則仍然可以發射，執行。（發射是順序的，執行時亂序的）
因此將指令的執行分為4級：
1。 發射：指令譯碼 並 檢測結構冒險（ID1） ，按照指令順序發射
指令的功能部件沒有結構競爭和沒有WAW冒險的時候，則這條指令可以發射。將會把指令發射到相應的功能部件，同時修改記分牌的內部資料結構。如果存在結構競爭或者是存在WAW冒險時候，指令暫停發射。
2。 讀運算元：等待到沒有資料冒險，再讀取操作（亂序讀） ，亂序讀運算元
記分牌監控源運算元是否就緒。一個源運算元就緒的條件為： 早前發射的活動指令對該運算元不進行寫入操作（即無RAW冒險）記分牌在這一步解決了RAW冒險問題。 當源運算元準備就緒，記分牌通知功能單元讀出運算元，並開始執行。 檢測RAW， 若有，則停頓該指令。但是在動態排程時，有多條指令並行操作，所以可能有另外指令滿足條件，則繼續執行下去，從而消除了停頓的損失。
3。 執行：對運算元進行操作 （EX），亂序執行
功能單元開始對運算元執行操作。當得到“結果”後，功能單元通知記分牌該操作已執行完畢。
4。 寫結果：完成執行 （WB），順序寫結果
檢查是否有WAR冒險，如果存在，則暫停指令。否則就寫入暫存器。 記分牌在各執行步驟中需檢測和記錄的事件。示例
DIVD F0，F2，F4 ADDD F10，F0，F8 SUBD F8，F8，F14
，記分板將暫停SUBD指令，直到ADDD指令讀取了運算元。
3、 GPU微架構模型
以下內容參考 1。 GPGPU-SIM Manual 2。 GPGPU-SIM-Presentation-On-Micro42
3。1 概覽 GPU
硬體體系結構
（假想模型）
GPGPU-Sim 所模擬的 GPU 結構如上圖所示。與現實中的 GPU 硬體結構相對應，其功能模組由三部分組成，分別為流多處理器（Stream Multiprocessors， SM）， 儲存器系統以及它們之間的網際網路絡。模擬器所模擬的每個 SIMT核心 中都包含有一個類似於簡單 MIPS 五段順序的流水線結構，SIMT核心簇透過虛擬的網際網路絡連線到儲存器子系統。每個儲存控制器控制兩個片外的 GDDR3/5
晶片
模型。最新版本的 GPGPU-Sim 模擬器添加了對現實中 GPU 執行緒處理簇（TPC，Thread Processing Cluster）的支援，這使得多個 SM 之間共用一個到網際網路絡的介面訪問資料。
3。2 SIMT 核心簇（假想模型）
如上圖所示，SIMT核心簇是包含一系列的SIMT核心，核心簇內部所有SIMT核心共享一個內部網際網路絡埠。核心簇共享一個FIFO佇列，用來儲存從網際網路絡取到的資料。這些個資料被定向傳送到SIMT核心的指令快取記憶體、或其儲存訪問單元 。為了使每個SIMT核心均能處理LD\ST指令，每個核心的LDST單元均有埠與外部請求埠相連、但同一簇內所有核心共用一個請求緩衝區。
3。3 GPU中SIMT核心（模擬對應實際的SM核心）
3。3。1 SIMT 核心概覽
每個SIMT核心模擬了一個SIMD處理器，其大致相當於nVidia所言的SM（Streaming Multiprocessor ，流式多處理器），或相當於AMD所言的CU （Compute Unit），SIMT核心的組織如上所示。
3。3。2 細探SIMT核心
SIMT核心被分為前端、後端，並且配有三個獨立的排程器。每部分
實際的SM對應於上圖的一個SIMT，而實際的SP核心對應於一組ALU流水線。
3。3。2。1 SIMT核心前端
FETCH 取指令（one instruction， one cycle， per warp）
這裡出現了第一個排程器：取指排程器。負責將取到的指令送入 I-Cache 中。
SIMT 前端中的指令快取模組（ I-Buffer）用於快取從 指令Cache中取出的指令。I-Buffer 被靜態劃分，使得執行在 SIMT 核心上的所有 warp 在其中都有專門的空間儲存指令。在當前的模型中，每個 warp 有兩個 I-Buffer 條目。 每個 I-Buffer 條目有一個有效位（valid bit），一個就緒位 （ready bit） 以及一個對應於該 warp 當前指令的譯碼後的指令。有效位表示當前 I-Buffer 的該條目中還有一個未發射的指令。 （該條目上的指令有效），就緒位表示該 Warp 當前指令已經準備好被髮射到執行流水線中。
通常情況，沒有結構冒險、沒有WAW衝突的時候，就緒位置1。
如果在 I-Buffer 中沒有任何有效指令，所有需要取指的 warp 會以輪詢的方式訪問 I-Cache 。一旦被選中，一個讀請求以及下一條指令的地址被送入到 I-Cache 。預設情況下，兩條連續的指令被取出。只要一個 Warp 被取指排程器排程進行取指操作，對應的 I-Buffer 條目有效位即為1， 直到該 Warp 內所有指令均執行完畢。
只有當一個執行緒執行完所有指令，並且沒有未完成寫回儲存器、寫回暫存器請求時，才能說一個執行緒執行完畢；只有當一個Warp內所有執行緒均執行完畢，該Warp才被認為執行完畢，並且不再受取值排程器排程；只有一個執行緒塊內所有Warp執行完畢，該執行緒塊才被認為執行完畢；只有所有塊執行完畢，該
核心函式函式
才算執行完畢。
DECODE 指令譯碼 （one instruction， one cycle， per warp）
在譯碼階段，當前被取出的指令被譯碼，確定指令種類（算術/分支/訪存）和要被使用的暫存器。之後便儲存到 I-Buffer 相應的條目中等待被髮射。譯碼的同時也會檢查暫存器的記分牌（score board），以確定可能有相關性的衝突。一旦檢測到衝突，將清空譯碼階段的輸入流水線暫存器，使正在譯碼的指令失效。若沒有衝突，將在記分牌入口設定標識，表示這些指令的流出的暫存器正被使用。
ISSUE 指令發射（multi instructions， one cycle， per warp）
這裡出現了第二個排程器：發射排程器，功能是從 I-Buffer 中選擇一個 Warp 發射到後續流水線中。此排程器獨立於之前的取指排程器。排程方式是迴圈優先順序策略（指不同Warp）。
發射排程器可以進行配置，每個週期從同一個 warp 中發射多條指令。更進一步地
GT200 （e。g。 Quadro FX 5800）： 允許雙Warp同時發射。
Fermi架構：Warp奇偶獨立排程器。
被髮射的指令必須滿足以下條件
該 warp 沒有處於柵欄同步（barrier）的等待狀態
（CUDA允許同一Block內不同Thread間實現塊內通訊）
I-Buffer 對應條目中的有效位為1（為0說明該條目上的指令無效）
透過記分板檢測
指令流水線
中的取運算元階段（operand access stage）不是掛起狀態。
發射指令目的地
儲存器相關指令（load， store ， memory barriers etc），被髮射到儲存流水線（參看下文 MEM PIPELINE）
運算指令被髮射到ALU計算單元，其包括多個SP流水線、SFU流水線。
分支指令將清空 I-Buffer 內所有與該 Warp 相關指令（參看下文SIMT STACK）
SIMT STACK：SIMT指令棧
每個 Warp 均有一個SIMT-stack，來解決一個Warp內的執行緒指令分歧問題。考慮到每個Warp內所有執行緒必須執行相同指令，因此，當不同執行緒出現不同分支情況時，所有執行緒將序列執行，這將會大大降低硬體的效率，因此我們需要一個方案，來降低這種分歧帶來的影響，最簡單的方法就是 PDOM（post-dominator stack-based reconvergence ） 機制。
以下參考 Dynamic Warp Formation： Efficient MIMD Control Flow on SIMD Graphics Hardware
上圖 a&b 兩個子圖說明了一個假想執行緒塊內四個執行緒的分支情況：其中 a 圖中每個 Flow Graph 指令中、按位代表該執行緒經過這個路徑。
精髓在於 c、d、e三個子圖：現在考慮四個執行緒，T1， T2， T3，T4 均執行到了 A指令 ，T1， T2， T3下一條指令為B，匯合點為G。（見 c 圖）
當執行完分支指令B時，Stack更新狀態成 d 圖，更新過程如下：
說明：TOS：Top Of Stack 棧首元素；x棧：某個 x 子圖所示棧；大寫字母，執行流上的指令；小寫字母：子圖
1）c棧的TOS，與d棧的TOS相同。只有Next PC 域被改變，變成了三個執行緒匯合點E；
2） B的一個分支 （D） 壓入棧（見標號d棧ii），連同將 D 分支對應的活動掩碼、
匯合點
E入棧，活動掩碼按位指示T2， T3執行緒執行該分支。
3）B的另一個分支（C）壓入棧（見標號d棧iii），連同將 C 分支對應的活動掩碼，匯合點E入棧，同樣的，活動掩碼按位指示 T1 執行緒執行該分支。
當執行B指令的後續分支指令時，開始彈棧：
1）當 T1 執行緒執行後續指令時，彈棧，棧頂元素的活動掩碼指示 T1 執行 C 分支，匯合點是 E （繼續彈棧，直到Next PC域為E）
2）當 T2， T3 執行緒執行後續指令時，彈棧，棧頂元素的活動掩碼指示 T2， T3 執行 D 分支，匯合點是 E （繼續彈棧，直到Next PC域為E）
3） 到達匯合點 E 時，SIMT stack如 e 所示，繼續重複上述過程。
nVidia 並沒有明示他們是如何處理分支指令的，在PTX檔案中也沒有有效資訊（見上述分析），但如果反編譯的話（cuobjdump），會得到和上述方法相同的結果。
SCOREBOARD 記分板
記分板部件檢查結構冒險、WAW衝突。如上所述，被某個Warp的一條指令寫入的暫存器在發射階段被預留。記分板裝置依靠Warp的ID進行索引。它儲存了對應某個Warp ID指令的所需要的暫存器的數目。預留的暫存器在寫回階段被釋放。
3。3。2。2 SIMT核心後端
OPRAND COLLECTOR 運算元收集器（本節用OP。COL。代替）
以下參考專利US7834881B2 - Operand collector architecture
上圖為假想圖，根據nVidia釋出的專利推測出來的Oprand Colloctor具體結構。
注意，這裡出現了第三個排程器，稱之為取數排程器。
組成 OP。COL。 的是一組
緩衝器
、和一個排程器。
每當一條指令被譯碼後，OP。COL。便為該指令分配空間，用於取數。OP。COL。 單元並沒有透過暫存器換名技術來消除暫存器名字依賴，而是透過另一種方式：確保每個週期內，對一個 Bank 的訪問，不得超過一次。
觀察上圖，其包含四個Collector Units，每個Unit包含三個運算元條目，和一個
標誌符
，用於指示當前該Unit屬於哪個Warp的哪條指令。
每個運算元條目包含四個欄位：
一個就緒位
一個有效位
一個暫存器識別符
運算元：該域包含128位元組，可以存放32個4位元組數，可以滿足一個Warp內的32個執行緒。
注意：每個Thread 有自己的暫存器，因此僅需一個暫存器標誌符即可
另外，排程器為每個Bank均保留了一個請求讀佇列，直到所有Unit對該Bank的訪問均已完畢。
當一條指令經過譯碼階段並且存在Collector Unit可用，則該Collector Unit 被分配給該指令，相應的Warp標誌、有效位、
暫存器識別
符被設定，運算元域被初始化為 0 。此外，運算元的讀請求被排隊到排程相應Bank佇列。實際上，執行單元寫回的資料的優先順序總是高於讀請求。排程器每週期選擇一組至多4個無bank衝突的資料傳送到暫存器堆。實際晶片中，為了減少 Crossbar 和 Collector Unit 的面積，每個Collector Unit 每個週期只接收一個運算元。
當每個運算元被從暫存器堆讀出並放入到相應的 OP。COL。 單元中，該指令就緒位被設定為1。最終，當一個指令的所有運算元都就緒後，該指令被髮射到一個SIMD執行單元。 實際上，對於每種不同的SIMD執行單元（SP，SFU，MEM），均有各自獨立的 Collector Units ，同時也有一組共享的 Collector Units。
MEM PIPELINE 儲存流水線
本階段處理執行緒訪問全域性記憶體與共享記憶體發出的請求。 每個 SIMT 核心有 4 種不同的片上一級儲存器：
共享儲存器
（ Shared Memory ）， 一級資料快取（ L1-data-cache ）， 常量快取 （ constant cache ）以及紋理快取（ Texture Cache ）。雖然上述四個儲存器物理上獨立，但由於其均為是儲存流水線 （LDST unit） 的組成部分，因此它們共享同一個寫回階段。
ALU UNIT 計算單元
GPGPU-Sim有兩種ALU計算單元：SP計算單元執行除超越方程外的任何指令；SFU計算單元執行超越方程額指令（sin，cos，log，etc）這兩種單元均以SIMD方式執行：
SP計算單元通常每週期執行一個Warp的一條指令
SFU計算單元執行週期視指令不同而不同：
sin指令需要4個週期
取倒數指令需要兩個週期
3。3。2。3 SIMT核心總覽
綜合前文分析，把每一部分串聯到一起，就得到了下圖。
