OpenCL:一種異構計算架構
轉載 王博部落格，https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
OpenCL:一種異構計算架構
目錄
1 摘要
由於電晶體功耗、物理效能的限制，CPU的發展受到了很大約束。 人們轉而尋找其它方式來提高系統性能，如多核處理器，異構平臺等。 開放式計算語言（OpenCL）的出現為當前大量存在的異構系統的平行計算提供了一個 標準。OpenCL通過一系列API的定義，提供硬體獨立的程式語言，為程式設計師提供 了靈活而又高效的程式設計環境。 本文通過對OpenCL計算架構的深入討論，指出了OpenCL程式設計的優勢及不足。並進行了 相關程式設計實踐，通過對不同裝置的並行程式設計測試，表明如果採用OpenCL並行程式設計架構， 能顯著提高程式的執行效率。
就目前的情況來看，異構系統有很高的價效比。相信在不久的將來，OpenCL將會成為 計算機並行、異構計算的重要組成部分。
關鍵字：OpenCL，異構計算，CPU/GPU計算，平行計算
2 為什麼需要OpenCL？
在過去的幾十年裡，計算機產業發生了巨大的變化。計算機效能 的不斷提高為當前各種應用提供了有力的保障。對於計算機的速度 而言，正如摩爾定律描述的那樣，是通過電晶體數目增加來提高頻率 的方式實現的。但是到了二十一世紀初期以後，這種增長方式受到 了一些限制，電晶體尺寸變得已經很小，其物理特性決定了很難再通 過大規模地增加電晶體的數目來提升頻率，且由於功耗也以非線性 的速度增加，因此這種方式受到很大的限制。在未來，這一趨勢會繼續 成為影響計算機系統最為重要的因素之一。
為了解決這一問題通常有兩種方式，第一種是通過 增加處理器的核心數目來為多工，多執行緒等提供支援，從整體 上提升系統的效能。第二種方式是通過異構的方式，例如可 利用CPU（Central Processing Unit）、GPU（Graphic Processing Unit）、甚至APU(Accelerated Processing Units， CPU與GPU的融合)等計算裝置的計算能力從而 來既提高系統的速度。
異構系統越來越普遍，對於支援這種環境的計算而言，也正受到越來越多 的關注。當前，不同廠商通常僅僅提供對於自己裝置程式設計的實現。對於異 構系統一般很難用同種風格的程式語言來實現機構程式設計，而且將不同的裝置 作為統一的計算單元來處理的難度也是非常大的。
開放式計算語言（Open Computing Language:OpenCL），旨在滿足這一重要需求。 通過定義一套機制，來實現硬體獨立的軟體開發環境。利用OpenCL可以充分利 用裝置的並行特性，支援不同級別的並行，並且能有效對映到由CPU，GPU， FPGA（Field－Programmable Gate Array）和將來出現的裝置 所組成的同構或異構，單裝置或多裝置的系統。OpenCL定義了執行時， 允許用來管理資源，將不同型別的硬體結合在同種執行環境中，並且很有希望 在不久的將來，以更加自然的方式支援動態地平衡計算，功耗和其他資源。
我相信在不久的將來，OpenCL將在異構並行程式設計中得到廣泛的應用。
3 OpenCL架構
3.1 介紹
OpenCL為異構平臺提供了一個編寫程式，尤其是並行程式的開放的框架標準。 OpenCL所支援的異構平臺可由多核CPU、GPU或其他型別的處理器組成。 OpenCL由兩部分組成，一是用於編寫核心程式(在OpenCL裝置上執行的程式碼) 的語 言，二是定義並控制平臺的API。OpenCL提供了基於任務和基於資料兩種並行計 算機制，它極大地擴充套件了GPU 的應用範圍，使之不再侷限於圖形領域。
OpenCL由Khronos Group維護。Khronos Group是一個非盈利性技術組織，維護著 多個開放的工業標準，例如OpenGL和OpenAL。這兩個標準分別用於三維圖形和計 算機音訊方面。
OpenCL源程式既可以在多核CPU上也可以在GPU上編譯執行，這大大提高了程式碼的 效能和可移植性。OpenCL標準由相應的標準委員會制訂，委員會的成員來自業界 各個重要廠商（主要有：AMD，Intel，IBM和NVIDIA）。作為使用者和程式設計師期待 已久的東西，OpenCL帶來兩個重要變化： 一個跨廠商的非專有軟體解決方案；一個跨平臺的異構框架以同時發揮系統中 所有計算單元的能力。
OpenCL支援廣泛的應用，將開發應用的過程一般化比較困難， 但是，通常來說，一個基於異構平臺的應用主要包含下面的 步驟[ 3 ]：
- 找出組成異構平臺的所有元件。
- 考察元件的特徵，這樣就能使得軟體根據不同的硬體特徵來實現。
- 建立在平臺上執行的一組核心。
- 設定與計算相關的儲存物件。
- 在合適的元件上以正確的順序執行核心。
- 收集結果。
這些步驟通過在OpenCL內部的一系列API和核心程式設計環境來實現。 這一實現採用“分治”策略。 可將問題分解為下面的模型[ 1 ] 平臺模型 執行模型 儲存模型 程式設計模型
這些概念是OpenCL整體架構的核心。 這四個模型將貫穿在整個OpenCL的程式設計過程中。
下面就簡要介紹這四個模型的相關內容。
3.2 平臺模型
平臺模型（如圖1）指定有一個處理器（主機Host）來協調程式的執行， 一個或多個處理器（裝置Devices）來執行OpenCL C程式碼。 在這裡其實僅僅是一種抽象的硬體模型，這樣就能方便程式設計師 編寫OpenCL C函式（稱之為核心）並在不同的裝置上執行。
圖中的裝置可以被看成是CPU/GPU，而裝置中的計算單元可以被看成是 CPU/GPU的核，計算單元的所有處理節點作為SIMD單元或SPMD單元（每個 處理節點維護自己的程式計數器）執行單個指令流。 抽象的平臺模型更與當前的GPU的架構接近。
平臺可被認為是不同廠商提供的OpenCL API的實現。如果一個平臺選定之後一般只能 執行該平臺所支援的裝置。就當前的情況來看，如果選擇了Intel的OpenCL SDK 就只能使用Intel的CPU來進行計算了，如果選擇AMD的APP SDK則能進行AMD的CPU和AMD的 GPU來進行計算。一般而言，A公司的平臺選定之後不能與B公司的平臺進行通訊。
3.3 執行模型
在執行模型中最重要的是核心，上下文和命令佇列的概念。上下文管理多個裝置， 每個裝置有一個命令佇列，主機程式將核心程式提交到不同的命令佇列上執行。
3.3.1 核心
核心是執行模型的核心，能在裝置上執行。當一個核心執行之前，需要指定一個 N-維的範圍（NDRange）。一個NDRange是一個一維、二維或三維的索引空間。 還需要指定全域性工作節點的數目，工作組中節點的數目。如圖NDRange所示， 全域性工作節點的範圍為{12, 12}，工作組的節點範圍為{4, 4}，總共有9個工作組。
例如一個向量相加的核心程式：
__kernel void VectorAdd(__global int *A, __global int *B, __global int *C){
int id = get_global_id(0);
C[id] = A[id] + B[id];
}
如果定義向量為1024維，特別地，我們可以定義全域性工作節點為1024， 工作組中節點為128，則總共有8個組。定義工作組主要是為有些僅需在 組內交換資料的程式提供方便。當然工作節點數目的多少要受到裝置的限制。 如果一個裝置有1024個處理節點，則1024維的向量，每個節點計算一次就能完成。 而如果一個裝置僅有128個處理節點，那麼每個節點需要計算8次。合理設定 節點數目，工作組數目能提高程式的並行度。
3.3.2 上下文
一個主機要使得核心執行在裝置上，必須要有一個上下文來與裝置進行互動。 一個上下文就是一個抽象的容器，管理在裝置上的記憶體物件，跟蹤在裝置上 建立的程式和核心。
3.3.3 命令佇列
主機程式使用命令佇列向裝置提交命令，一個裝置有一個命令佇列，且與上下文 相關。命令佇列對在裝置上執行的命令進行排程。這些命令在主機程式和裝置上 非同步執行。執行時，命令間的關係有兩種模式：(1)順序執行，(2)亂序執行。
核心的執行和提交給一個佇列的記憶體命令會生成事件物件。 這用來控制命令的執行、協調宿主機和裝置的執行。
3.4 記憶體模型
一般而言，不同的平臺之間有不同的儲存系統。例如，CPU有快取記憶體而GPU就沒有。 為了程式的可移植性，OpenCL定義了抽象的記憶體模型，程式實現的時候只需關注抽 象的記憶體模型，具體向硬體上的對映由驅動來完成。記憶體空間的定義及與硬體的映 射大致如圖所示。
記憶體空間在程式內部可以用關鍵字的方式指定，不同的定義與資料存在的位置 相關，主要有如下幾個基本概念[ 2 ]：
- 全域性記憶體：所有工作組中的所有工作項都可以對其進行讀寫。工作項可以 讀寫此中記憶體物件的任意元素。對全域性記憶體的讀寫可能會被快取，這取決於裝置 的能力。
- 不變記憶體：全域性記憶體中的一塊區域，在核心的執行過程中保持不變。 宿主機負責對此中記憶體物件的分配和初始化。
- 區域性記憶體：隸屬於一個工作組的記憶體區域。它可以用來分配一些變數， 這些變數由此工作組中的所有工作項共享。在OpenCL裝置上，可能會將 其實現成一塊專有的記憶體區域，也可能將其對映到全域性記憶體中。
- 私有記憶體：隸屬於一個工作項的記憶體區域。 一個工作項的私有記憶體中所定義的變數對另外一個工作項來說是不可見的。
3.5 程式設計模型
OpenCL支援資料並行，任務並行程式設計，同時支援兩種模式的混合。對於同步 OpenCL支援同一工作組內工作項的同步和命令佇列中處於同一個上下文中的 命令的同步。
4 基於OpenCL的程式設計示例
在本小節中以影象旋轉的例項，具體介紹OpenCL程式設計的步驟。 首先給出實現流程，然後給出實現影象旋轉的C迴圈實現和OpenCL C kernel實現。
4.1 流程
4.2 影象旋轉
4.2.1 影象旋轉原理
影象旋轉是指把定義的影象繞某一點以逆時針或順時針方向旋轉一定的角度， 通常是指繞影象的中心以逆時針方向旋轉。假設影象的左上角為(l, t), 右下角為(r, b)，則影象上任意點(x, y) 繞其中心(xcenter, ycenter)逆時針旋轉θ角度後， 新的座標位置(x',y')的計算公式為：
x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,
y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.
C程式碼：
void rotate(
unsigned char* inbuf,
unsigned char* outbuf,
int w, int h,
float sinTheta,
float cosTheta)
{
int i, j;
int xc = w/2;
int yc = h/2;
for(i = 0; i < h; i++)
{
for(j=0; j< w; j++)
{
int xpos = (j-xc)*cosTheta - (i - yc) * sinTheta + xc;
int ypos = (j-xc)*sinTheta + (i - yc) * cosTheta + yc;
if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
outbuf[ypos*w + xpos] = inbuf[i*w+j];
}
}
}
OpenCL C kernel程式碼：
#pragma OPENCL EXTENSION cl_amd_printf : enable
__kernel void image_rotate(
__global uchar * src_data,
__global uchar * dest_data, //Data in global memory
int W, int H, //Image Dimensions
float sinTheta, float cosTheta ) //Rotation Parameters
{
const int ix = get_global_id(0);
const int iy = get_global_id(1);
int xc = W/2;
int yc = H/2;
int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;
int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;
if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H))
dest_data[ypos*W+xpos]= src_data[iy*W+ix];
}
旋轉45度
正如上面程式碼中所給出的那樣，在C程式碼中需要兩重迴圈來計算橫縱座標上新的 座標位置。其實，在影象旋轉的演算法中每個點的計算可以獨立進行，與其它點的 座標位置沒有關係，所以並行處理較為方便。OpenCL C kernel程式碼中用了並行 處理。
上面的程式碼在Intel的OpenCL平臺上進行了測試，處理器為雙核處理器，影象大小 為4288*3216，如果用迴圈的方式執行時間穩定在0.256s左右，而如果用OpenCL C kernel並行的方式，執行時間穩定在0.132秒左右。GPU的測試在NVIDIA的GeForce G105M顯示卡 上進行，執行時間穩定在0.0810s左右。從迴圈的方式，雙核CPU並行以及GPU平行計算 已經可以看出，OpenCL程式設計的確能大大提高執行效率。
5 總結
通過對OpenCL程式設計的分析和實驗可以得出，用OpenCL編寫的應用具有很好的移 植性，能在不同的裝置上執行。OpenCL C kernel一般用並行的方式處理，所以能極大地提高程式的執行效率。
異構平行計算變得越來越普遍，然而對於現今存在的OpenCL版本來說，的確還存在 很多不足，例如編寫核心，需要對問題的並行情況做較為深入的分析，對於記憶體的 管理還是需要程式設計師來顯式地申明、顯式地在主存和裝置的儲存器之間進行移動， 還不能完全交給系統自動完成。從這些方面，OpenCL的確還需加強，要使得人們能高 效而又靈活地開發應用，還有很多工作要完成。
6 參考文獻
【1】 Aaftab Munshi. The OpenCL Specification Version1.1 Document Revision:44[M]. Khronos OpenCL Working Group. 2011.6.1.
【2】Aaftab Munshi. 倪慶亮譯. OpenCL規範 Version1.0 Document Revision:48[M]. Khronos OpenCL Working Group. 2009.10.6.
【3】Aaftab Munshi, Benedict R. Gaster, Timothy G. Mattson, James Fung, Dan Ginsburg. OpenCL Programming Guide [M]. Addison-Wesley Professional. 2011.7.23.
【4】Benedict Gaster, Lee Howes, David R. Kaeli and Perhaad Mistry. Heterogeneous Computing with OpenCL[M]. Morgan Kaufmann, 1 edition. 2011.8.31.
【5】Slo-Li Chu, Chih-Chieh Hsiao. OpenCL: Make Ubiquitous Supercomputing Possible[J]. IEEE International Conference on High Performance Computing and Communications. 2010 12th 556-561.
【6】John E. Stone, David Gohara, Guochun Shi. OpenCL: A parallel programming standard for heterogeneous computing systems[J]. Copublished by the IEEE CS and the AIP. 2010.5/6 66-72.
【7】Kyle Spafford, Jeremy Meredith, Jeffrey Vetter. Maestro:Data Orchestration and Tuning for OpenCL Devices[J]. P. D'Ambra,M.Guarracino, and D.Talia (Eds.):Euro-Par 2010,Part II,LNCS6272, pp. 275–286, 2010. \copyright Springer-Verlag Berlin Heidelberg 2010.