292      技術社區[雲棲]

從零開始學習OpenCL開發（二）一個最簡單的示例與簡單性能分析

歡迎關注 轉載請注明 https://blog.csdn.net/leonwei/article/details/8893796

1 Hello OpenCL 

    這裏編寫一個最簡單的示例程序，演示OpenCl的基本使用方法：

    1.首先可以從Nvdia或者Amd或者Intel或者所有OpenCl成員的開發者網站上下載一份他們實現的OpenCL的SDK。雖然不同公司支持了不同版本的OpenCL和擴展ext，但是在相同版本上對於標準的OpenCL接口，每個SDK實現的結果都是一樣的，如果你隻是用標準的OpenCL規範，那麼采用哪個SDK無所謂，當然有些公司把OpenCL SDK捆綁在更大的SDK裏，如NVDIA放在他們的CUDA開發包裏，這時我們要做的隻是把其中cl文件夾下的h 以及 OpenCL.lib OpenCL.dll文件拿出來就行。

   下麵進入代碼的部分，本例中實現兩個一維數組的相加（這是最容易理解的可並行計算問題），代碼主要這幾個部分：

   2.獲取機器中所有已實現的OpenCL平台：

   //get platform numbers
   err = clGetPlatformIDs(0, 0, &num);
 

 //get all platforms
  vector<cl_platform_id> platforms(num);
  err = clGetPlatformIDs(num, &platforms[0], &num);

  首先要知道OpenCL平台platform是什麼意思。我們知道不同OpenCL組織裏不同廠商的不同硬件都紛紛支持OpenCL標準，而每個支持者都會獨自去實現OpenCl的具體實現，這樣如果你的機器中有很多個不同“OpenCl廠商”的硬件（通常實現在驅動中），那麼你的機器中就會出現幾套對OpenCL的不同實現，如你裝了intel cpu，可能就一套intel的實現，裝了NVDIA的顯卡，可能還有一套Nvidia的實現，還有值得注意的是，就算你可能沒有裝AMD的顯卡，但是你裝了AMD的opencl開發包，你機器中也可能存在一套AMD的實現。這裏的每套實現都是一個platform，可以說不同廠商拿到的SDK可能是一樣的，但是查詢到的機器裏的platform則可能是不一樣的，sdk是代碼層，platform是在驅動裏的實現層，opencl在不同廠商的代碼層一樣，但是在一個機器裏會存在不同的實現層（原涼我這麼囉嗦，但是這個問題我開始糾結了很久）。

 不同廠商給了相同的代碼SDK，但是在驅動層，不同廠商的實現是完全不一樣的，也就是paltform是不一樣的，例如NVIDIA的的platform隻支持N自己的顯卡作為計算設備（可能他們認為cpu作為計算設備是在是雞肋），但是AMD的platform則不僅支持AMD自己的設備，還支持Intel的CPU。

所以你要在程序開始查詢機器所有支持的platform，再根據情況選擇一個合適的paltform。（通常你要選擇包含compute device的能力最強的那個platform，例如你發現客戶機裝的是N卡，而機器上有N的platform那麼就選它了）

通過clGetPlatformInfo 這個函數還可以進一步的得到該平台的更多信息（名字、cl版本、實現者等等）

  3.查詢device信息（在程序中這一步是可以不做的，但是可以用來判斷platform的計算能力）

  //get device num

 err=clGetDeviceIDs(platforms[0],CL_DEVICE_TYPE_ALL,0,0,&num);
  vector<cl_device_id> did(num);

 //get all device
  err=clGetDeviceIDs(platforms[0],CL_DEVICE_TYPE_ALL,num,&did[0],&num);

  //get device info

 clGetDeviceInfo（...）

 以上代碼可以獲取某個platform下的所有支持的device（這裏和下麵都特指compute device，因為在pc下host device一定是你的CPU了）

 這些有助於你判斷用哪個platform的計算能力更強

 4.選定一個platform，創建context（設備上下文）

 cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };

 cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_ALL, NULL, NULL, &err);

 上麵代碼首先使用你選定的那個paltform設置context屬性，然後利用這個屬性創建context。context被成功創建好之後，你的CL工作環境就等於被搭建出來了，CL_DEVICE_TYPE_ALL意味著你把這個platform下所有支持的設備都連接進入這個context作為compute device。

 5.為每個device創建commandQueue。command queue是像每個device發送指令的信使。

   cqueue[i] = clCreateCommandQueue(context, did[0], 0, 0);

6.下麵進入真正在device run code的階段：kernal函數的準備

    首先準備你的kernal code，如果有過shader編程經驗的人可能會比較熟悉，這裏麵你需要把在每個compute item上run的那個函數寫成一段二進製字符串，通常我們實現方法是寫成單獨的一個文件（擴展名隨意），然後在程序中使用的時候二進製讀入這個文件。

   例如本例的數組相加的kernal code：

 __kernel void adder(__global const float* a, __global const float* b, __global float* result)
{
 int idx = get_global_id(0);
 result[idx] = a[idx]) +b[idx];
}

具體的限定符和函數我們後麵會分析，但是這段代碼的大意是獲取當前compute item的索引idx，然後兩個數組idx上的成員相加後存儲在一個buf上。這段代碼會盡可能並行的在device上跑。

把上麵那個文件命名為kernal1.cl

然後在程序中讀入它到字符串中（通常你可以為這個步驟寫一個工具函數）

ifstream in(_T("kernal11.cl"), std::ios_base::binary);
 if(!in.good()) {
  return 0;
 }

 // get file length
 in.seekg(0, std::ios_base::end);
 size_t length = in.tellg();
 in.seekg(0, std::ios_base::beg);

 // read program source
 std::vector<char> data(length + 1);
 in.read(&data[0], length);
 data[length] = 0;

 // create and build program
 const char* source = &data[0];

這樣我們的kernal code就裝進char* source裏麵了。

7.從kernal code 到program

 program在cl中代表了程序中所用到的所有kernal函數及其使用的函數，是device上代碼的抽象表示，我們需要把上麵的char* source轉化成program：

cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);

clBuildProgram(program, 0, 0, 0, 0, 0)

如上兩句代碼分別先從字符串的source創建一個program，在build它（我們說過OpenCl是一個動態編譯的架構）

8 .  拿到kernal 函數

kernal是CL中對執行在一個最小粒度的compute item上的代碼及參數的抽象（你可以理解成為cpu上的main函數）。

我們需要首先從前麵build好的program裏抽取我們要run的那個kernal函數。

cl_kernel adder = clCreateKernel(program, "adder", 0);

9. 準備kernal函數的參數

kernal函數需要三個參數，分別是輸入的兩個數組mem，和一個輸出的數組mem，這些mem都要一一創建準備好。

首先是輸入的兩個mem

std::vector<float> a(DATA_SIZE), b(DATA_SIZE)
 for(int i = 0; i < DATA_SIZE; i++) {
  a[i] = i;
  b[i] = i;

 }

a個b是我們要運算的兩個輸入數組（注意他們是在CPU上的，或者說分配與你的主板內存）

cl計算的變量要位於device的存儲上（例如顯卡的顯存），這樣才能快起來，所以首先要把內存搬家，把這部分輸入數據從host mem拷貝到device的mem上，代碼如下：

cl_mem cl_a = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &a[0], NULL);
 cl_mem cl_b = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(cl_float) * DATA_SIZE, &b[0], NULL);

上麵代碼的含義是使用host mem的指針來創建device的隻讀mem。

最後還要在device上分配保存結果的mem

cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);

這是直接在device上分配的。

最後設置好kernal的參數

 clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);
 clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);
 clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);

10.執行kernal函數

err = clEnqueueNDRangeKernel(cqueue[0], adder, 1, 0, &work_size, 0, 0, 0, 0);

注意cl的kernal函數的執行是異步的，這也是為了能讓cpu可以與gpu同時做事（但是異步就涉及到設備間的同步、狀態查詢等，這是非常複雜的一部分，後麵再說）

所以上麵這個函數會立即返回，clEnqueueNDRangeKernel的意思是往某個device的commoand queue裏麵推入一個kernal函數讓其執行，device會按某個順序執行它的command queue裏麵的指令，所以這個語句調用後，kernal是否真的立即執行還要取決於它的queue裏麵是否還有其他的指令。

11.將結果拷回CPU

  上麵執行後的結果是直接寫在device的存儲上，通常要在代碼中繼續使用，我們就需要把這個結果再拷回到CPU的內存上，使用下麵的代碼：

std::vector<float> res(DATA_SIZE)

 err = clEnqueueReadBuffer(cqueue[0], cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0);

clEnqueueReadBuffer的含義是往command queue裏麵推出一個條指令，是回拷mem，這裏麵的CL_TRUE是標誌著這個指令的執行的同步的，就會阻塞cpu，所以這行代碼返回就標誌著該device上直到這個指令之前的所有指令都已經執行完了。

 上麵為止就可以到帶在res裏我們使用cl在device上執行kernla函數的結果了，可以與純CPU的執行結果對比一遍，結果應該是一致的。

12.打掃戰場

//release
 clReleaseKernel(adder);
 clReleaseProgram(program);
 clReleaseMemObject(cl_a);
 clReleaseMemObject(cl_b);
 clReleaseMemObject(cl_res);

 for(size_t i=0;i<num;i++){
  clReleaseCommandQueue(cqueue[i]);
 }
 clReleaseContext(context);

2.性能分析

 上麵的是一個非常簡單的CL入門程序。借助這個程序，我後來又做了很多性能分析，想知道究竟使用CL執行運算和平常的CPu上運算有什麼區別，性能會有怎樣的不同。

我修改了不同版本的kernal函數，使kernal的運算複雜度不斷提升，並在不同platform下和單純在CPU上執行這些運算，得到的統計數據如下：

注意：

0.1、2、3的複雜度分別使用的簡單擴大數組長度、求冪操作、增加求冪操作的指數

1.以下的數據皆為毫秒

2.第一列為傳統的CPU運算，後兩列為使用Amd 和Nvidia兩個平台的運算

3.由於測試機未安裝AMD顯卡，所以AMD平台使用的device其實是一個CPU，所以1、2、3列代表的情況可以看做純CPU，使用openCL架構用CPU做計算設備、使用OpenCL架構用GPU做設備

4.由於OpenCL架構多涉及到一個host和device間內存拷貝的操作，2、3列中的+號兩端分別代表拷貝內存所用的時間和實際運算時間。

從上表我們“以偏蓋全”的得到一些結論：

1.純CPU的計算會隨著計算複雜度的增加而顯著上漲，純GPU的CL架構的計算在與此同時計算耗時基本平穩，雖然在第一個運算，GPU的時間還會高於CPU，但是到第三個運算時GPU的時間依然沒有明顯增長，而CPU已經長到GPU時間的70多倍。

2.不同平台的CL實現在內存拷貝上所化時間基本一致，這部分時間跟計算複雜度無官，隻跟內存大小有關。在我們的例子中他們都是63ms

3.從1.2列的對比看出，就算是同樣使用CPU做為計算，在CL架構下性能也會得到較大提升，雖然實質上1和2列都是最終在CPU上計算，但是CL的架構可能封轉了更高一層，利用了CPU內的一些高級指令或者利用了CPU的更多的並行計算能力。

4.OpenCL是真正兼容各種硬件的，不同於CUDA，這對於產業化產品的開發意義重大，在主流的機器上，你總能找到一個可用的opencl platform，而它都會比CPU計算提示性能。

   從這個簡單的性能分析可以看出，使用OpenCL架構的異構計算可以大幅度提高傳統在CPU上的計算性能，而且這種提高可能會隨著計算量的複雜度升高而增長，所以那些所謂“百倍”、“千倍”的增長在某些計算領域是有可能的，同時盡量使用GPU做device是可以最大提升性能的；

  同時我們要注意到異構計算通常涉及到大量的內存拷貝時間，這取決於你內存與顯存間的帶寬，這部分時間是不可忽視的，如果一個計算工作，它在CPU上運行的時間都比內存在異構設備間拷貝的時間短，那麼將他做OpenCL的加速是沒有任何意義的，也就是說我們要注意計算的複雜度，複雜度過小的計算使用異構計算反而會增加計算時間，GPU運算都存在一個跟計算複雜度無關的“起步時間”（例如本例在180ms左右，當計算在CPU上執行小於180ms時放在GPU上是無意義的。）

最後更新：2017-04-03 18:51:45

  上一篇： 2013第四屆藍橋杯 C/C++本科A組 真題答案解析【交流帖】

  下一篇： Android在自定義View（SurfaceView）中實現進度條Progress