simt編譯

1. mali是手機端gpu系列嗎

簡單介紹下Arm Mali的GPU系列

對於手機終端來說，GPU圖像處理能力是衡量一台手機的性能標桿。首先，是UI流暢性，大家拿到手機都得先劃來劃去看下UI是否流暢，而UI其實主要還是用GPU渲染的；其次是游戲的流暢性，對於很酷炫的游戲，GPU是目前手機端的唯一高性能3D加速器。在手機端，主流的幾個GPU主要是PowerVr，Mali，Adreno。蘋果早起使用的就是PowerVr的定製版本，不過隨著蘋果自研GPU，PowerVr現在基本可以是算賣給了紫光；Mali是鼎兄並鼎大名的安謀半導體ARM的圖形加速IP；Adreno是高通的圖形GPU。當然這里不是要對比這些GPU的性能，而是簡單介紹下Mali的GPU系列。

Mali其實是Arm的Mali系列IP核，但是很多現在在很多網上提到Mali其實是直接認為是Mali的GPU。Mali系列其實還有視頻，顯示控制器，camera等。但是Mali應該算是授權比較多的。而且因為GPU也被更多的非業內人士所熟知。

1、Mali的四大架構之一：Utgard

第一代微架構Utgard（北歐神話人物：烏特加德）。這一代架構出來的比較早，主要是圖形加速IP。可以追溯到2007年的mali-200。不過最讓人驚訝的是mali-4xx系列，現在很多電視晶元都還在用這個IP。比如小米的智能電視，還有很多是mali-4xx系列的。

Utgard這一代vertex shader和fragment shader是分離的，arm官方支持的Opengl ES也只維護到2.0。所以Opengl ES 3.0及其以上要求的app是跑不了的。並且OpenCL在這一代中也是不支持的，因為這一代主打圖形計算，不怎麼支持通用計算。

移動端的GPU主要以基於tile的計算為主岩旁，mali的全系列（截止目前）都是基於tile的計算。基於tile的計算可以簡單的認為，在進行計算的時候一張圖會被劃分成若干張小矩形，每個矩形可以認為是一個tile，然後對一個tile進行同時計算。

主要系列有：mali-200, mali-400, mali-450, mali-470

2、Mali的四大架構之二：Midgard

第二代微架構Midgard（北歐神話人物：米德加德）。Midgard這一代GPU開始屬於同一著色器的架構，也就是上面說的vertex shader和fragment shader已經統一在一起了，粗塵橡相當於同一個shader計算單元可以處理多種著色器。當然也開始支持計算通用計算。特別是對OpenCL的支持，對通用計算有了很大的支持。OpenGLES 3.1雖然引入了compute shader，但是說起通用計算，OpenCL顯然更加專業。

這個架構是基於128bit向量的，所以在編程的時候往往用4個float編程了能最大發揮其性能。當然，編譯器也會把某些可以進行優化的計算合並成向量進行計算，不過最好在編碼階段自行優化。編譯器編譯的優化比較難以去把握。當然，也不建議用大於128bit的方式進行編程，最終需要編譯器拆成多個數的運算，且每個數的位寬最大為128bit，如果編譯器優化不好，反而會導致性能下降。

主要系列有：mali-t6xx, mali-t7xx, mali-t8xx

3、Mali的四大架構之三：Bifrost

第三代微架構Bifrost（北歐神話中連接天宮和大地的：彩虹橋）。由於這一代產品基本在2016年後發布的了，而OpenGLES在2016年後基本穩定了，所以相對於Midgard來說，在大方向上圖形計算這塊也沒有多大的需要調整。

在Bifrost（Bifrost上更像是SIMT的模式，這里用SIMT表述也是我從多個文檔資料推敲出來的）上會先把向量拆成標量，然後每個線程跑多維向量的第一維，因此對於三維向量 vec3向量最快只需要3個cycle，對於思維向量vec4最快只需要4個cycle。這里用了最快這個表述是因為並不是所有的指令都是單個cycle的。

當然，雖然bifrost架構是標量運算的，這是針對32bit的位寬來說的，如果是16bit位寬的計算，一個線程是可以在一個cycle內處理一個vec2的16bit數據的。因此在編程的時候，如果是8bit或者16bit的數據，用於應該考慮如何組織代碼使得更有效的組合運算，例如16bit位寬的情況，盡量是用vec2，8bit位寬的盡量用vec4。

對於Bifrost，例如G76，一個shader core可以同時運行幾十個線程,，從mali的資料顯示，shader core一般由三個部分組成，ALU,L/S,TEXTURE三個主要模塊。在G76上是8-wide wrap的，一般設置為3個ALU。（其餘的型號可能不一樣，例如G51/G72是4-wide wrap的，G72同樣是3個ALU；G52跟G76一樣，不過G52可配置成2個ALU的）對於AI加速方面，部門系列也有一些指令修改，例如G52和G76都引入了int8 dot指令，該指令針對神經卷積網路的運算做了優化。

主要系列有：mali-g31, mali-g51, mali-g71, mali-g52, mali-g72, mali-g76

4、Mali的四大架構之四：Valhall

第四代微架構Valhall是2019年第二季度推出來的。該系列的是基於超標量實現的。對於G77，使用的時16-wide的wrap，單個shader core集成兩個計算引擎。

主要系列有：mali-g57, mali-g77

最後，本文簡要的梳理了下mali gpu架構的一些情況，同時對不同架構上的一些計算資源進行簡要描述，希望能給看到的朋友提供一些有用的信息。同時，也吐槽一下，Mali系列的晶元命名在Bifrost和Valhal系列沒有區分開，單純從mali-g51, mali-g71, mali-g52, mali-g72，mali-g57, mali-g77，很難區分最後兩個型號是Valhall的架構。這個命名不知道mali是怎麼考慮的，著實令人難解。

2. CUDA能取代CPU嗎

摩爾的預言 唯有CUDA才是終極的CPU 架構解析：從硬體角度看CUDA（上）如果你是一位小熊在線的忠實讀者，那麼你對NVIDIA最新的GPU架構技術一定並不陌生。如果你不經常看小熊在線的文章，那麼我還是建議你多做一些功課，去了解一下NVIDIA最新的顯卡特色技術。因為NVIDIA的CUDA架構與他們的GPU架構稍有不同，並且在這篇文章中，將為讀者揭示一些NVIDIA從未提及的技術細節。GPU的處理核心架構從上面的這張架構圖我們可以看到，NVIDIA的Shader核心是由若干個材質處理單元（TPC）的群組所構成的。例如一個8800GTX，它內部就有8個群組，而8800GTS內部就有6個群組。事實上，每一個群組都是由1個材質處理單元和2個流多重處理器構成的。而處理器又由一個前端的讀取/解碼單元，一個指令發送單元，一個由八個計算單元組成的組，和2個SFU超級功能單元所組成，他們在處理指令方面都屬於SIMD單指令多數據流。同樣這類指令也適用於所有warp中的線程。NVIDIA這種並行模式叫做SIMT單指令多線程執行單元。需要指出的是，後端的操作，比前端操作要高出2倍的時鍾周期。在實際中，部分執行指令有可能有2倍長。16路的SIMD，就只能相當於8路的處理效能。流多重處理器在運作模式下，允許每個時鍾周期內，在流水線的前端，讓一個warp處在准備執行的狀態。在一個warp中的32個線程中的應用指令，在流水線的後端，需要用4個時鍾周期來處理。但是這種操作在流水線的前端，僅僅需要1/2的時鍾周期就可以完成，從這一點來看，僅僅需要2個時鍾周期，就可以完成操作。 因此為了保證流水線前後端的速度一致，使得硬體的利拆圓頌用率最大化，CUDA推出了一種交替類型的指令周期。每個典型的指令使用一個時鍾周期，而SFU的指令在另一個周期。每個旅鄭多重處理器也有一些確定數量的資源，以便可以更精確的榨取它們的性能。它們有一些很小的存儲區域，這些區域被叫做共享內存。每一個多重處理器的共享內存僅有16KB大小。這並不是一個高速緩存，這是由程序自由管理的存儲區域。這方面與Cell處理器上的SPU處理單元有些類似，可以進行本地的數據存儲。這些技術細節是非常獨特而有腔中趣的，並且你可以在CUDA的文檔中找到相關更詳細的軟體與硬體技術說明。這些存儲區域並不是為了pixel shaders而設計的，他們主要針對的就是CUDA架構解析：從硬體角度看CUDA（下）這個存儲區域為block塊中的線程提供了一種信息溝通的途徑。它最重要的作用就是強調限制級別。所有的線程都保證被封裝到一個block塊中，這就保證了多重處理器可以有效的處理這些任務。反過來說，指派block塊到不同的多重處理器中，是非常不確定的。這就意味著，在執行的時候來自不同block塊中的線程是無法進行通信的，使用這種存儲區是一件非常復雜的事兒。但是設計這些復雜的存儲區，對於整體架構來說也是值得的。因為某些特殊的線程，可能會貿然訪問主內存，這也許會產生很多沖突，而這種共享存儲區可以快速的鏈接寄存器。block塊內部架構這些共享的存儲區並不是多重處理器唯一可以訪問的存儲設備。很顯然，他們可以使用顯存，但是相對於共享存儲區來說，顯存的帶寬和速度都不如前者。因此，這種機制可以抑制對內存過於頻繁的訪問。NVIDIA也提供了多重處理器的高速緩存，可以存儲常量和紋理，大致相當於每個多重處理器能分配到8KB的空間。多重處理器內部架構同時，多重處理器也具備8192個寄存器。在多重處理器中所有激活的block塊中所有的線程，都可以通過這些寄存器共享信息。不過激活的warps總數被限制在24個，也就是768個線程。例如8800GTX最多可在同一時間處理12288個線程。在這方面加以限制，是為了不讓GPU的資源消耗的太多，可以更合理的分配計算任務。針對CUDA優化過的程序從本質上講就是在block塊的數量與他們的尺寸之間找到一種平衡。在一個block塊中加入更多的線程，有利於提高內存潛伏期的效率，但是與此同時可以使用的寄存器數量就要少了。如果block塊中的線程數太多了，比如達到512個線程的水平，那麼整個流水線的執行效能就大大降低了。這僅僅夠喂飽一個多重處理器的，浪費了256個線程的處理能力。因此NVIDIA建議每個block塊使用128至256個線程，這是最為折中的辦法，它會在內核的潛伏期與寄存器數量之間找到最好的平衡點架構解析：從軟體角度看CUDA從軟體的角度來看，CUDA是由C語言的擴展組成，這種架構當然要追溯到BrookGPU的時代，並且許多API的編寫工作，都是從BrookGPU繼承來的。CUDA擴展和延伸了許多BrookGPU的函數和變數。這里有一個非常典型的關鍵字__global__，它是一個前綴，在它後面的內容意味著是一個內核。這個功能將被CPU和執行單元所調用。__device__，標記有這個關鍵詞的程序將會被GPU執行。但是僅僅在GPU中被調用。最後再介紹一個關鍵詞__host__，這個前綴是可選的，它會指派CPU調用一個函數，並且由CPU執行它，換言之，它就像是一個非常傳統的CPU執行函數。對於__device__ 和 __global__ 函數來說，也有一些限制條件。他們不能被遞歸，也就是說他們不能自己調用自己。他們不能擁有可變數目的自變數。最後，要說一下__device__關鍵詞，它必須駐留在GPU的顯存空間內。理論上來說，絕對不可能獲得它的物理地址。變數在控制存儲區域的時候，也會有一些限制。__shared__ 是一個變數的前綴，當程序中出現這個關鍵詞的時候，就表明變數要保存在流多重處理器中的共享存儲區中。這與通過 __global__關鍵詞調用會有些不同。這是因為執行配置中規定，在調用時，要指明grid柵格容器的大小，並指明它在哪個內核中，並且要為每一個block塊指派所包含線程的尺寸。我們可以看看下面的這個例子：其中Dg是grid柵格的尺寸，而Db是定義block塊，這兩個變數是CUDA的一種新的類型。CUDA API其本質上來講是由各種操作顯存的函數組成的。cudaMalloc用來分配內存，cudaFree用來釋放內存，cudaMemcpy用來互相拷貝內存和顯存之間的數據。在最後，我們將介紹一下CUDA程序的編譯方式。這是非常有趣的，整個編譯過程需要幾個階段。首先，所有的代碼都要讓CPU來處理，這些都會從文件中提取，並且他們都會通過標準的編譯器。用於GPU處理的代碼，首先要轉換成中間媒介性語言——PTX。中間語言更像是一種匯編程序，並且能夠中和潛在的無效代碼。在最後的階段，中間語言會轉換成指令。這些指令會被GPU所認同，並且會以二進制的形式被執行。

3. GTX690的詳細信息

流處理器暴增之謎
基於效能和計算能力方面的考慮，NVIDIA與AMD不約而同的改變了架構，NVIDIA雖然還是採用SIMT架構，但也借鑒了AMD「較老」的SIMD架構之作法，降低控制邏輯單元和指令發射器的比例，用較少的邏輯單元去控制更多的CUDA核心。於是一組SM當中容納了192個核心的壯舉就變成了現實！
通過右面這個示意圖就看的很清楚了，CUDA核心的縮小主要歸功於28nm工藝的使用，而如此之多的CUDA核心，與之搭配的控制邏輯單元面積反而縮小了，NVIDIA強化運算單元削減控制單元的意圖就很明顯了。
此時相信有人會問，降低控制單元的比例那是不是意味著NVIDIA賴以成名的高效率架構將會一去不復返了？理論上來說效率肯定會有損失，但實際上並沒有想像中的那麼嚴重。NVIDIA發現線程的調度有一定的規律性，編譯器所發出的條件指令可以被預測到，此前這部分工作是由專門的硬體單元來完成的，如今可以用簡單的程序來取代，這樣就能節約不少的晶體管。
所以在開普勒中NVIDIA將一大部分指令派發和控制的操作交給了軟體（驅動）來處理。而且GPU的架構並沒有本質上的改變，只是結構和規模以及控制方式發生了變化，只要驅動支持到位，與游戲開發商保持緊密的合作，效率損失必然會降到最低——事實上NVIDIA著名的The Way策略就是干這一行的！
這方面NVIDIA與AMD的思路和目的是相同的，但最終體現在架構上還是有所區別。NVIDIA的架構被稱為SIMT（Single Instruction Multiple Threads，單指令多線程），NVIDIA並不像AMD那樣把多少個運算單元捆綁為一組，而是以線程為單位自由分配，控制邏輯單元會根據線程的任務量和SM內部CUDA運算單元的負載來決定調動多少個CUDA核心進行計算，這一過程完全是動態的。
但不可忽視的是，軟體預解碼雖然大大節約了GPU的晶體管開銷，讓流處理器數量和運算能力大增，但對驅動和游戲優化提出了更高的要求，滾清這種情況伴隨著AMD度過了好多年，NVIDIA也要面對相同的問題了，希望他能做得更好一些，否則散熱量還將繼續增加。 SMX與SM的改動細節
全新的Kepler相比上代的Fermi架構改變了什麼，看架構圖就很清楚了：
GK104相比GF110，整體架構沒有大的改變，GPU（圖形處理器集群）維持4個，顯存控制器從6個64bit（384bit）減至4個64bit（256bit），匯流排介面升級至PCIE 3.0。剩下的就是SM方面的改變了
NVIDIA把GK104的SM（不可分割的流處理器集群）稱為SMX，原因就是暴增的CUDA核心數量。但實際上其結構與上代的SM沒有本質區別，不同的只是各部分單元的數量和比例而已飢悉。具體的區別逐個列出來進行對比： Kepler與Fermi架構SM參數對比 單元 GF100 GF104 GK104 GK104/GF104 CUDA 32 48 192 4：1 SFU 4 8 32 4：1 Warp 2 2 4 2：1 Dispatch 2 4 8 2：1 LD/ST 16 16 32 2：1 TMU 4 8 32 4：1 1. NVIDIA把流處理器稱為CUDA核心；
2. SFU（Special Function Units，特殊功能單元）是比CUDA核心更強的額外運算單元，可用於執行抽象的指令，例如正弦、餘弦、倒數和平方根，圖形插值指令也在SFU上執行；
3. Warp是並行線程調度器，每一個Warp都可以調度SM內部的所有CUDA核心或者SFU；
4. Dispatch Unit是指令分派單元，分則將Warp線程中的指令按照順序和相關性分配給不同的CUDA核心或SFU處理；
5. LD/ST就是載入/存儲單元，可以為每個線程存儲運算源地址與路徑，方便隨時隨地的從緩存或顯存中存取數據；
6. TMU是紋理單元，用來處理紋理和陰影貼圖、屏幕空間環大肢前境光遮蔽等圖形後期處理；
通過以上數據對比不難看出，GK104暴力增加CUDA核心數量的同時，SFU和TMU這兩個與圖形或計算息息相關處理單元也同比增加，但是指令分配單元和線程調度器還有載入/存儲單元的佔比都減半了。這也就是前文中提到過的削減邏輯控制單元的策略，此時如何保證把指令和線程填滿一個CUDA核心，將是一個難題。據知名評測網站卡吧基地最新資訊顯示，此難題將在8個月內得到初步解決。

4. GPU硬體基礎知識

GPU channel 是GPU與CPU之間的橋接介面，通過CPU向GPU發送GPU指令的唯一通道，GPU channel包含了兩類用於存儲GPU指令的buffer：

當GPU指令被寫入到GPU command buffer時，系統還會向Ring buffer中寫入與此指令所對應的packet，packet包含了此指令在GPU command buffer中的偏移位置與長度數據。

在執行指令的時候，GPU不是直接從GPU command buffer中讀取數據，而是先經過Ring buffer讀取出當前待處理指令的相關信息，再據此讀取GPU command（這也是為什麼Ring buffer被稱之為indirect buffer的原因）。

現代GPU為了加強數據的並行化處理的強度，使用的是SIMT（Single Instruction Multi Thread，SIMD的更高級版本）體系結構，shader program運行的最小單位是thread，多個運行相同shader的threads會被打包到一個組（此滾thread group），這個thread group，在NVIDIA被稱之為warp，在AMD中被稱之為wavefront。

上面這張圖是從標題鏈接給出的Turing白皮書中截取的GPU架構圖，其中包含如下幾個關鍵縮寫：

GPU中用於存儲數據的結構有多種[4]，分別是：

每種存儲結構都有著各自的優缺點此山，因此適用於不同的應用場景，從訪問速度來看，這些存儲結構按照從高到低排序依次是：
RMEM > SMEM > CMEM > TMEM > LMEM > GMEM

RMEM與SMEM是直接集成在GPU晶元上的，而剩下的幾種存儲結構則是在GPU之外的晶元上的，此外，LMEM/CMEM/TMEM都有著各自的緩存機制，即在訪問數據的時候都會首先從緩存中進行查找判斷，再決定是否需要從更低一級速度的存儲結構中進行讀取。

存儲在LMEM中的數據可見性與RMEM一樣，都是只對負責對其進行讀寫的線程可見。LMEM實際上森扒中並不是一塊物理存儲空間，而是對GMEM的一個抽象，因此其訪問速度與對GMEM的訪問速度是相同的。LMEM中的數據對於一個線程而言是Local的（即只從屬於當前thread的空間，對其他線程不可見），通常用於存儲一些automatic變數（automatic變數指的是一些大尺寸的數據結構或者數組，因為寄存器不夠，因此會塞入LMEM中），編譯器在寄存器不足的時候，就會從GMEM中開辟一塊空間用作LMEM。

雖然LMEM是從GMEM中分割出來的，但是其使用方式與GMEM還是有著一些區別：

如上圖所示（從圖中可以看出，L1是位於GPU晶元上的，其中SMEM就存儲在其中，RMEM也是在晶元上，而L2及以後的存儲空間則都是晶元之外的存儲空間了），在對LMEM進行數據讀寫的時候，會經歷這樣一個緩存層級流動：L1->L2->LMEM。因為LMEM實際上是臨時開辟的一塊空間，因此裡面的數據實際上是GPU先寫入的，在此之前發生的讀取就相當於讀到了一堆亂碼。

那麼什麼情況下會使用到LMEM呢？一般來說有如下兩種情形：

因為LMEM相對於寄存器訪問速度的低效性，因此其對性能的影響主要有如下兩個方面：

但是因為以下的兩點原因，LMEM也不一定會造成性能下降：

對於一些LMEM可能會存在瓶頸的情況，參考文獻[3]中給出了一些分析的方法可供排查，同時還給出了對應的優化策略以及實戰案例，有興趣的同學可以前往參考。

存儲在RMEM中的數據只對負責對此寄存器進行讀寫的線程可見，且其生命周期與此線程的生命周期一致。

通常情況下，對寄存器的訪問不需要消耗時鍾周期，但是在一些特殊情況（比如先進行了一個寫操作，之後再進行讀取，或者在bank訪問沖突的情況下），會有例外。先寫後讀的延遲大概是24個時鍾周期，對於更新的GPU（每個SM包含32個cores的情況），可能需要花費768個線程來隱藏這個延遲。

當需求的寄存器數目超出硬體所能支持的限額時，就會導致寄存器壓力，在這種情況下，數據就會使用LMEM來進行存儲（所謂的spilled over，即溢出），如下圖所示[3]：

存儲在SMEM中的數據對處於同一個block所有的線程都是可見的（不負shared之名），因此通常用於多個線程之間的數據互換，為了避免多個線程同時訪問相同的數據導致的阻塞，NVIDIA將SMEM劃分成32個邏輯單元，每個單元叫做一個bank，在內存中連續的數據，在banks的分布也是連續的：

SMEM是位於L1 Cache中的，其尺寸通常為16/32/48KB，剩餘部分用作L1 Cache，對於開普勒架構而言，每個bank每個時鍾的帶寬是64bits/clock，較早的Fermi架構時鍾不太一樣，但是帶寬差不多是這個數值的一半。

由於一個warp中有32個線程，因此總共需要32個SMEM banks。由於每個bank在每個時鍾周期中只支持一次訪問請求，因此多個同時訪問的請求就會導致bank conflict，這個的處理過程後面會講。

默認每個bank佔用32bits（4bytes），開普勒架構之後，可以通過指令（cudaDeviceSetSharedMemConfig()）將每個bank擴充到64bits，以應對雙精度數據的訪問沖突。

存儲在Global Memory中的數據對於當前進程中的所有線程都是可見的，其生命周期與進程一致。

CMEM通常用於存儲一些常量數據，當同一個warp中的所有線程都需要使用同一個參數時，可以將數據放在CMEM中，這種做法比將數據放在GMEM中更節省帶寬。

TMEM也是一種常量存儲結構，當一個warp中的線程所需要讀取的數據都是存儲位置上相鄰的時候，使用這種結構比GMEM具有更優的性能表現（也是出於帶寬的原因）

[1]. A HISTORY OF NVIDIA STREAM MULTIPROCESSOR
[2]. Life of a triangle - NVIDIA's logical pipeline
[3]. Local Memory and Register Spilling
[4]. GPU Memory Types – Performance Comparison

5. 計算機組成與設計：硬體/軟體介面的目錄

出版者的話
譯者序
前言
第1章計算機概要與技術1
1.1引言1
1.1.1計算應用的分類及其特性2
1.1.2你能從本書學到什麼3
1.2程序概念入門4
1.3硬體概念入門7
1.3.1剖析滑鼠8
1.3.2顯示器8
1.3.3打開機箱9
1.3.4數據安全12
1.3.5與其他計算機通信13
1.3.6處理器和存儲器製造技術14
1.4性能15
1.4.1性能的定義15
1.4.2性能的測量17
1.4.3CPU性能及其因素18
1.4.4指令的性能19
1.4.5經典的CPU性能公式19
1.5功耗牆21
1.6滄海巨變：從單處理器向多處理器轉變23
1.7實例：製造以及AMD Opteron X4基準25
1.7.1SPEC CPU基準測試程序27
1.7.2SPEC功耗基準測試程序28
1.8謬誤與陷阱29
1.9本章小結31
1.10拓展閱讀32
1.11練習題32
第2章指令：計算機的語言42
2.1引言42
2.2計算機硬體的操作43
2.3計算機硬體的操作數46
2.3.1存儲器操作數47
2.3.2常數或立即數操作數49
2.4有符號和無符號數50
2.5計算機中指令的表示54
2.6邏輯操作59
2.7決策指令61
2.7.1循環62
2.7.2case/switch語句64
2.8計算機硬體對過程的支持65
2.8.1使用更多的寄存器66
2.8.2嵌套過程68
2.8.3在棧中為新數據分配空間69
2.8.4在堆中為新數據分配空間70
2.9人機交互72
2.10MIPS中32位立即數和地址的定址75
2.10.132位立即數75
2.10.2分支和跳轉中的定址76
2.10.3MIPS定址模式總結78
2.10.4機器語言解碼79
2.11並行與指令：同步81
2.12翻譯並執行程序83
2.12.1編譯器84
2.12.2匯編器84
2.12.3鏈接器85
2.12.4載入器87
2.12.5動態鏈接庫87
2.12.6啟動一個Java程序88
2.13以一個C排序程序為例89
2.13.1swap過程89
2.13.2sort過程90
2.14數組與指針95
2.14.1用數組實現clear96
2.14.2用指針實現clear96
2.14.3比較兩個版本的clear97
2.15高級內容：編譯C語言和解釋Java語言98
2.16實例：ARM指令集98
2.16.1定址模式99
2.16.2比較和條件分支100
2.16.3ARM的特色100
2.17實例：x86指令集101
2.17.1Intel x86的改進101
2.17.2x86寄存器和數據定址模式103
2.17.3x86整數操作104
2.17.4x86指令編碼106
2.17.5x86總結107
2.18謬誤與陷阱107
2.19本章小結108
2.20拓展閱讀110
2.21練習題110
第3章計算機的算術運算135
3.1引言135
3.2加法和減法135
3.2.1多媒體算術運算137
3.2.2小結138
3.3乘法139
3.3.1順序的乘法演算法和硬體139
3.3.2有符號乘法141
3.3.3更快速的乘法142
3.3.4MIPS中的乘法142
3.3.5小結142
3.4除法143
3.4.1除法演算法及其硬體結構143
3.4.2有符號除法145
3.4.3更快速的除法146
3.4.4MIPS中的除法146
3.4.5小結147
3.5浮點運算148
3.5.1浮點表示149
3.5.2浮點加法152
3.5.3浮點乘法154
3.5.4MIPS中的浮點指令157
3.5.5算術精確性162
3.5.6小結164
3.6並行性和計算機算術：結合律165
3.7實例：x86的浮點165
3.7.1x86浮點體系結構166
3.7.2Intel SIMD流擴展2（SSE2）浮點體系結構167
3.8謬誤與陷阱168
3.9本章小結170
3.10拓展閱讀172
3.11練習題173
第4章處理器182
4.1引言182
4.1.1一個基本的MIPS實現183
4.1.2實現方式概述183
4.2邏輯設計慣例185
4.3建立數據通路187
4.4一個簡單的實現機制192
4.4.1ALU控制192
4.4.2主控制單元的設計194
4.4.3數據通路的操作197
4.4.4控制的結束199
4.4.5為什麼不使用單周期實現方式201
4.5流水線概述202
4.5.1面向流水線的指令集設計205
4.5.2流水線冒險205
4.5.3對流水線概述的小結210
4.6流水線數據通路及其控制211
4.6.1圖形化表示的流水線219
4.6.2流水線控制222
4.7數據冒險：轉發與阻塞225
4.8控制冒險234
4.8.1假定分支不發生234
4.8.2縮短分支的延遲235
4.8.3動態分支預測237
4.8.4流水線小結239
4.9異常240
4.9.1異常在MIPS體系結構中的處理241
4.9.2在流水線實現中的異常242
4.10並行和高級指令級並行245
4.10.1推測的概念246
4.10.2靜態多發射處理器247
4.10.3動態多發射處理器250
4.11實例：AMD Opteron X4（Barcelona）流水線253
4.12高級主題：通過硬體設計語言描述和建模流水線來介紹數字設計以及更多流水線示例255
4.13謬誤與陷阱255
4.14本章小結256
4.15拓展閱讀257
4.16練習題257
第5章大容量和高速度：開發存儲器層次結構280
5.1引言280
5.2cache的基本原理283
5.2.1cache訪問285
5.2.2cache缺失處理288
5.2.3寫操作處理289
5.2.4一個cache的例子:內置FastMATH處理器290
5.2.5設計支持cache的存儲系統292
5.2.6小結294
5.3cache性能的評估和改進295
5.3.1通過更靈活地放置塊來減少cache缺失297
5.3.2在cache中查找一個塊300
5.3.3替換塊的選擇302
5.3.4使用多級cache結構減少缺失代價302
5.3.5小結305
5.4虛擬存儲器305
5.4.1頁的存放和查找308
5.4.2缺頁309
5.4.3關於寫312
5.4.4加快地址轉換：TLB312
5.4.5集成虛擬存儲器、TLB和cache315
5.4.6虛擬存儲器中的保護317
5.4.7處理TLB缺失和缺頁318
5.4.8小結322
5.5存儲器層次結構的一般架構323
5.5.1問題1：一個塊可以被放在何處323
5.5.2問題2：如何找到一個塊324
5.5.3問題3：當cache缺失時替換哪一塊325
5.5.4問題4：寫操作如何處理325
5.5.53C：一種理解存儲器層次結構行為的直觀模型326
5.6虛擬機328
5.6.1虛擬機監視器的必備條件329
5.6.2指令集系統結構（缺乏）對虛擬機的支持329
5.6.3保護和指令集系統結構329
5.7使用有限狀態機來控制簡單的cache330
5.7.1一個簡單的cache330
5.7.2有限狀態機331
5.7.3一個簡單的cache控制器的有限狀態機333
5.8並行與存儲器層次結構：cache一致性334
5.8.1實現一致性的基本方案335
5.8.2監聽協議335
5.9高級內容：實現cache控制器336
5.10實例：AMD Opteron X4(Barcelona)和Intel Nehalem的存儲器層次結構337
5.10.1Nehalem和Opteron的存儲器層次結構337
5.10.2減少缺失代價的技術339
5.11謬誤和陷阱340
5.12本章小結342
5.13拓展閱讀343
5.14練習題343
第6章存儲器和其他I/O主題355
6.1引言355
6.2可信度、可靠性和可用性357
6.3磁碟存儲器359
6.4快閃式存儲器362
6.5連接處理器、內存以及I/O設備363
6.5.1互聯基礎364
6.5.2x86處理器的I/O互聯365
6.6為處理器、內存和操作系統提供I/O設備介面366
6.6.1給I/O設備發送指令367
6.6.2與處理器通信368
6.6.3中斷優先順序369
6.6.4在設備與內存之間傳輸數據370
6.6.5直接存儲器訪問和內存系統371
6.7I/O性能度量：磁碟和文件系統的例子372
6.7.1事務處理I/O基準程序372
6.7.2文件系統和Web I/O的基準程序373
6.8設計I/O系統373
6.9並行性與I/O：廉價磁碟冗餘陣列374
6.9.1無冗餘(RAID 0)376
6.9.2鏡像（RAID 1）376
6.9.3錯誤檢測和糾錯碼（RAID 2）376
6.9.4位交叉奇偶校驗（RAID 3）376
6.9.5塊交叉奇偶校驗（RAID 4）376
6.9.6分布式塊交叉奇偶校驗（RAID 5）377
6.9.7P+Q冗餘（RAID 6）378
6.9.8RAID小結378
6.10實例：Sun Fire x4150伺服器379
6.11高級主題：網路383
6.12謬誤與陷阱383
6.13本章小結386
6.14拓展閱讀387
6.15練習題387
第7章多核、多處理器和集群394
7.1引言394
7.2創建並行處理程序的難點396
7.3共享存儲多處理器398
7.4集群和其他消息傳遞多處理器400
7.5硬體多線程403
7.6SISD、MIMD、SIMD、SPMD和向量機404
7.6.1在x86中的SIMD：多媒體擴展405
7.6.2向量機406
7.6.3向量與標量的對比407
7.6.4向量與多媒體擴展的對比408
7.7圖形處理單元簡介408
7.7.1NVIDIA GPU體系結構簡介410
7.7.2深入理解GPU411
7.8多處理器網路拓撲簡介412
7.9多處理器基準測試程序415
7.10Roofline：一個簡單的性能模型417
7.10.1Roofline模型418
7.10.2兩代Opteron的比較419
7.11實例：使用屋頂線模型評估四種多核處理器422
7.11.14個多核系統422
7.11.2稀疏矩陣424
7.11.3結構化網格425
7.11.4生產率426
7.12謬誤與陷阱427
7.13本章小結428
7.14拓展閱讀429
7.15練習題429
附錄A圖形和計算GPU439
A.1引言439
A.1.1GPU發展簡史439
A.1.2異構系統440
A.1.3GPU發展成了可擴展的並行處理器440
A.1.4為什麼使用CUDA和GPU計算440
A.1.5GPU統一了圖形和計算441
A.1.6GPU可視化計算的應用441
A.2GPU系統架構441
A.2.1異構CPU-GPU系統架構442
A.2.2GPU介面和驅動443
A.2.3圖形邏輯流水線443
A.2.4將圖形流水線映射到統一的GPU處理器443
A.2.5基本的統一GPU結構444
A.3可編程GPU445
A.3.1為實時圖形編程446
A.3.2邏輯圖形流水線446
A.3.3圖形渲染程序447
A.3.4像素渲染示例447
A.3.5並行計算應用編程448
A.3.6使用CUDA進行可擴展並行編程449
A.3.7一些限制453
A.3.8體系結構隱含的問題453
A.4多線程的多處理器架構454
A.4.1大規模多線程454
A.4.2多處理器體系結構455
A.4.3單指令多線程（SIMT）456
A.4.4SIMT warp執行和分支457
A.4.5管理線程和線程塊457
A.4.6線程指令458
A.4.7指令集架構（ISA）458
A.4.8流處理器（SP）461
A.4.9特殊功能單元（SFU）461
A.4.10與其他多處理器的比較461
A.4.11多線程多處理器總結462
A.5並行存儲系統462
A.5.1DRAM的考慮462
A.5.2cache463
A.5.3MMU463
A.5.4存儲器空間463
A.5.5全局存儲器463
A.5.6共享存儲器464
A.5.7局部存儲器464
A.5.8常量存儲器464
A.5.9紋理存儲器464
A.5.10表面465
A.5.11load/store訪問465
A.5.12ROP465
A.6浮點算術465
A.6.1支持的格式465
A.6.2基本算術465
A.6.3專用算術466
A.6.4性能467
A.6.5雙精度467
A.7資料：NVIDIA GeForce 8800468
A.7.1流處理器陣列（SPA）468
A.7.2紋理/處理器簇（TPC）469
A.7.3流多處理器（SM）470
A.7.4指令集471
A.7.5流處理器（SP）471
A.7.6特殊功能單元（SFU）471
A.7.7光柵化471
A.7.8光柵操作處理器（ROP）和存儲系統471
A.7.9可擴展性472
A.7.10性能472
A.7.11密集線性代數性能472
A.7.12FFT性能473
A.7.13排序性能474
A.8資料：將應用映射到GPU474
A.8.1稀疏矩陣475
A.8.2在共享存儲器中進行緩存477
A.8.3掃描和歸約478
A.8.4基數排序480
A.8.5GPU上的N-Body應用482
A.9謬誤與陷阱486
A.10小結489
A.11拓展閱讀489
附錄B匯編器、鏈接器和SPIM模擬器490
B.1引言490
B.1.1什麼時候使用匯編語言493
B.1.2匯編語言的缺點493
B.2匯編器494
B.2.1目標文件的格式495
B.2.2附加工具496
B.3鏈接器498
B.4載入499
B.5內存的使用499
B.6過程調用規范500
B.6.1過程調用502
B.6.2過程調用舉例503
B.6.3另外一個過程調用的例子505
B.7異常和中斷507
B.8輸入和輸出509
B.9SPIM511
B.10MIPS R2000匯編語言513
B.10.1定址方式514
B.10.2匯編語法515
B.10.3MIPS指令編碼515
B.10.4指令格式516
B.10.5常數操作指令520
B.10.6比較指令520
B.10.7分支指令521
B.10.8跳轉指令523
B.10.9陷阱指令523
B.10.10取數指令525
B.10.11保存指令526
B.10.12數據傳送指令527
B.10.13浮點運算指令528
B.10.14異常和中斷指令532
B.11小結533
B.12參考文獻533
B.13練習題533