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.

OpenCL:一種異構計算架構