參考書：《計算機體系結構量化研究方法》 作者：John L. Hennessy

一、 引言

指令流與數據流的并行分類

• 單指令流、單數據流（SISD）
• 單指令流、多數據流（SIMD）
• 多指令流、單數據流（MISD）基本不使用這種類型
• 多指令流、多數據流（MIMD）

SIMD相比于MIMD的優勢
由于數據操作是并行的，所以程序員可以采用順序思維方式但卻能獲得并行加速比

SIMD的三種變體

• 向量體系結構
• 多媒體SIMD指令集擴展
• 圖形處理單元（GPU）

二、 向量體系結構

本質：以流水線形式來執行多數據操作

2.1 VMIPS指令集體系結構

基于MIPS指令集，我們看一個向量處理器，VMIPS；它的標量部分為MIPS，它的向量部分是MIPS的邏輯向量擴展。主要組件如下：

示例：實現DAXPY（double precision ax plus y,ax+y）
MIPS代碼如下，

L.D FO,aDADDIU R4,Rx,#512 ;載入最后地址 LOOP: L.D F2,0(Rx) ；載入x[i]MUL.D F2,F2,F0L.D F4,0(Ry) ADD.D F4,F4,F2S.D F4,9(Ry)DADDIU` Rx,Rx,#8 ;遞增XDADDIU Ry,Ry,#8DSUBU R20,R4,Rx ;計算范圍BNEZ R20,LOOP ;檢查是否完成

DAXPY的VMIPS代碼：

L.D F0,aLV V1,RxMULVS.D V2,V1,F0LV V3,RyADDVV.D V4,V2,V2SV V4,Ry

向量處理器最大的好處就是大幅度縮短了動態指令帶寬，僅執行6條指令，而MIPS幾乎要執行600條。這一縮減是因為向量運算是對64個元素執行的，在MIPS中差不多占據一般循環的開銷指令在VMIPS中是不存在的。如果循環的迭代之間沒有相關性(循環間相關)，那么這些循環就可以向量化。

2.2 向量執行時間

幾個基礎概念

• 護航指令組：其中指令不能包含任何結構性冒險的一組向量指令
• 鐘鳴：用于估計護航指令組執行時間的度量單位，執行m個護航指令組構成的向量序列需要m次鐘鳴
• 鏈接：將不存在結構性冒險的數據鏈接（轉發）在一起

影響執行時間的因素

• 操作數向量的長度
• 操作之間的結構性冒險
• 數據相關

向量執行時間開銷源

• 發射限制
• 向量啟動時間

2.3 多車道

向量指令集的優點與特性

• 允許軟件使用一條很短的指令就能向硬件傳送大量并行任

務

• 所有向量算數指令只允許一條指令的向量寄存器N與其他向量寄存器的元素N進行運算

因此采用多車道來提高向量單元的峰值吞吐量。結構如圖：

從單車道變為了四車道后，就會將1次鐘鳴的時鐘周期數由64變為了16個。由于多車道非常有利，所以應用程序和體系結構都必須支持長向量；否則，它們的快速執行速度會耗盡指令帶寬。
需要注意的是四條車道彼此之間并不聯系，這樣避免通信就減少了構建高并行執行單元所需要的連接成本與寄存器端口。

2.4 向量長度寄存器：處理不等于64的循環

實際上，特定向量運算的長度在編譯時通常是未知的。如下代碼：
for (i=0;i<n;i++)
y[i] = a * x[i] + y[i];
所有的向量長度都取決于n,而它的取值在執行前不可能知道。
解決方案：向量長度寄存器（VLR）
但VLR中的不能超過最大向量長度（MLR）,但是一不小心大于MLR怎么辦，畢竟長度都是未知的。所以就需要采用條帶挖掘的技術
條帶挖掘：生成一些代碼，使每個向量運算都是針對小于或等于MVL的大小來完成的。一般條帶挖掘采用兩個循環：

• 循環處理迭代數為MVL倍數的情況
• 處理所有其他迭代及小于MVL的情況

如下是C語言中DAXPY條帶挖掘的版=本：

low = 0; VL = (n % MVL); //使用求模運算找出不規則大小部分 for (j=0;j<=(n/MVL);j++) {for (i=low;i<(low+VL);i++) //執行長度為VLy[i] = a * x[i] + y[i];low = low + VL;VL = MUL; //將長度復位為最大向量長度 }

主要思想是執行(n/MVL+1)次循環，在循環的第一次執行長向量的非規則部分。

2.5 向量遮罩寄存器：處理向量循環中的if語句

對于中低向量化級別的程序，加速比是非常有限的。導致循環向量化程度較低的兩個主要原因：

• 循環內部存在條件（if語句）
• 稀疏矩陣的存在

因此引入遮罩寄存器。遮罩寄存器可以用來實現一條向量指令中每個元素的條件執行。在啟用向量遮罩寄存器時，任何向量指令都只會針對符合特定條件的向量元素來執行，即這些元素在向量遮罩寄存器中的相應項目為1.如:
SNEVS.D V1,F0 ;只有V1向量不為F0，才繼續執行后續指令
SUBVV.D V1,F1,V2 ;在向量遮罩下執行減法
遮罩寄存器的開銷：

• 條件執行的指令在不滿足條件時也需要執行時間

三、 圖形處理器

GPU的祖先是圖形加速器，極強的圖形處理能力是GPU得以存在的原因。

GPU幾乎擁有所有可以由編程環境捕獲的并行類型：多線程、MIMD、SIMD、指令集并行。
NVIDIA開發了一種與C類似的語言和編程環境，通過客服異質計算及多種并行帶來的雙重挑戰來提高GPU程序員的生產效率。這一效率的名稱為CUDA：compute unified device architecture,計算統一設備體系。所有這些并行形式的統一主題就是CUDA線程。
以最低級別的并行作為編程原型，編譯器和硬件可以將數以千計的CUDA線程聚合在一起，利用GPU中的各種并行類型：多線程，MIMD、SIMD、指令級并行。因此NVIDIA將CUDA編程定義為：單指令多線程（SIMT）。這些線程進行了分塊，在執行時以32個線程為1組，稱為線程塊，我們馬上就會明白其原因。我們將執行整個線程塊的硬件稱為多線程SIMD處理器。
我們只需要幾個細節就能給出CUDA編程的示例。
為了區分GPU（設備）的功能與系統處理器（主機）的功能，CUDA使用__device__或__global__表示前者，使用__host__表示后者。
被聲明為__device__或__global__functions的CUDA變量被分配給GPU存儲器，可以供所有多線程SIMD處理器訪問。
對于在GPU上運行的函數name進行擴展函數調用的語法為：

name<<<dimGrid,dimBlock>>>(…parameter list…)

其中dimGrid和dimBlock規定了代碼的大小（用塊表示）和塊的大小（用線程表示）
除了塊識別符（blockIdx）和每個線程識別符（threadIdx）之外，CUDA還為每個塊的線程數提供了一個關鍵字（blockDim），它來自上一個細節中提到的dimBlock參數。一下是DAXPY的cuda代碼:

//調用DAXPY，每個線程塊中有256個線程 __host__ int nblocks = (n+255)/256;daxpy<<<nblocks,256>>>(n,2.0,x,y); //CUDA中的DAXPY __device__ void daxpy(int n,double a,double *x, double *y) {int i = blockIdx.x * blockDim.x +threadIdx.x;if (i<n) y[i] = a*x[i] + y[i]; }

如何理解tid為blockIdx.x * blockDim.x +threadIdx.x呢？
首先要了解網格與線程的關系，如下：

總結

以上是生活随笔為你收集整理的数据级并行--计算机体系结构的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。