[转]OpenCL 教学（一）
Contents
penCL 簡介
OpenCL是由Khronos Group針對異質性計算裝置（heterogeneous device）進行平行化運算所設計的標準API以及程式語言。所謂的「異質性計算裝置」，是指在同一個電腦系統(tǒng)中，有兩種以上架構差異很大的計算裝置，例如一般的CPU以及顯示晶片，或是像CELL的PPE以及SPE。目前，最為常見的就是所謂的GPGPU應用，也就是利用一般的顯示晶片（即GPU）進行3D繪圖以外的計算工作。過去GPGPU的應用，有各種不同的使用方式。最早的GPGPU，多半是直接透過3D繪圖的API進行，例如OpenGL或D3D的HLSL（High Level Shading Language）。但是，這樣做有很多缺點，主要是即使想要進行的運算和3D繪圖無關，仍需要處理很多3D繪圖方面的動作（例如建立texture，安排render-to-texture動作等等）。這讓GPGPU變得十分復雜。后來開始有些嘗試把這些3D繪圖部份隱藏起來的想法，例如由Stanford大學設計的BrookGPU，可以透過不同的backend將Brook程式轉換成由CPU、Direct3D、或OpenGL來執(zhí)行。另外，也有各家顯示卡廠商自行開發(fā)的系統(tǒng)，包括ATI針對其產(chǎn)品設計的Close to Metal（以及后來的AMD Stream），以及NVIDIA的CUDA。Microsoft也在DirectX 11中加入了特別為GPGPU設計的DirectCompute。
由于各家廠商的GPGPU 方案都是互不相容的（例如AMD Stream 的程式無法在NVIDIA 的顯示晶片上執(zhí)行，而CUDA 的程式也不能在AMD 的顯示晶片上執(zhí)行），這對GPGPU 的發(fā)展是不利的，因為程式開發(fā)者必須為不同廠商的顯示晶片分別撰寫程式，或是選擇只支援某個顯示晶片廠商。由于顯示晶片的發(fā)展愈來愈彈性化，GPGPU 的應用范圍也增加，因此Apple 決定提出一個統(tǒng)一的GPGPU 方案。這個方案得到包括AMD、IBM、Intel、NVIDIA 等相關廠商的支持，并很快就交由Khronos Group 進行標準化。整個計畫只花了五個月的時間，并在2008 年十二月時正式公開。第一個正式支援OpenCL 的作業(yè)系統(tǒng)是Apple 的MacOS X 10.6 "Snow Leopard"。AMD 和NVIDIA 也隨后推出了在Windows 及Linux 上的OpenCL 實作。IBM 也推出了支援CELL 的OpenCL 實作。
OpenCL 的主要設計目的，是要提供一個容易使用、且適用于各種不同裝置的平行化計算平臺。因此，它提供了兩種平行化的模式，包括task parallel 以及data parallel。目前GPGPU 的應用，主要是以data parallel 為主，這里也是以這個部份為主要重點。所謂的data parallel，指的是有大量的資料，都進行同樣的處理。這種形式的平行化，在很多工作上都可以見到。例如，影像處理的程式，經(jīng)常要對一個影像的每個pixel 進行同樣的動作（例如Gaussian blur）。因此，這類工作很適合data parallel 的模式。
OpenCL 的架構
OpenCL 包括一組API 和一個程式語言。基本上，程式透過OpenCL API 取得OpenCL 裝置（例如顯示晶片）的相關資料，并將要在裝置上執(zhí)行的程式（使用OpenCL 程式語言撰寫）編繹成適當?shù)母袷?#xff0c;在裝置上執(zhí)行。OpenCL API 也提供許多裝置控制方面的動作，例如在OpenCL 裝置上取得一塊記憶體、把資料從主記憶體復制到OpenCL 裝置上（或從OpenCL 裝置上復制到主記憶體中）、取得裝置動作的資訊（例如上一個程式執(zhí)行所花費的時間）等等。
例如，我們先考慮一個簡單的工作：把一群數(shù)字相加。在一般的C 程式中，可能是如下：
float a[DATA_SIZE];
float b[DATA_SIZE];
float result[DATA_SIZE];
// ...
for(int i = 0; i < DATA_SIZE; i++) {
result[i] = a[i] + b[i];
}在OpenCL 中，則大致的流程是：透過data parallel 的模式，這里的OpenCL 程式非常簡單，如下所示：
__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];
}在一般的版本中，是透過一個回圈，執(zhí)行DATA_SIZE次數(shù)的加法動作。而在OpenCL中，則是建立DATA_SIZE個work item，每個work item都執(zhí)行上面所示的kernel。可以看到，OpenCL程式語言和一般的C語言非常類似。__kernel表示這個函式是在OpenCL裝置上執(zhí)行的。__global則表示這個指標是在global memory中（即OpenCL裝置上的主要記憶體）。而get_global_id(0)會傳回work item的編號，例如，如果有1024個work item，則編號會分別是0 ~ 1023（實際上編號可以是二維或三維，但在這里先只考慮一維的情形）。
要如何讓上面這個簡單的OpenCL kernel 實際在OpenCL 裝置上執(zhí)行呢？這就需要透過OpenCL API 的幫助了。以下會一步一步說明使用OpenCL API 的方法。
OpenCL 環(huán)境設定
在使用OpenCL API 之前，不免要進行一些環(huán)境的設定。相關的動作可以參考下列的文章：
- 在Windows 下使用OpenCL
- 在Xcode 中使用OpenCL
開始撰寫OpenCL 程式
在使用OpenCL API之前，和絕大部份所有其它的API一樣，都需要include相關的header檔案。由于在MacOS X 10.6下OpenCL的header檔案命名方式和在其它作業(yè)系統(tǒng)下不同，因此，通常要使用一個#ifdef來進行區(qū)分。如下所示：
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
這樣就可以在MacOS X 10.6 下，以及其它的作業(yè)系統(tǒng)下，都可以include 正確的OpenCL header 檔。接著，要先取得系統(tǒng)上所有的OpenCL platform。在MacOS X 10.6 下，目前只有一個由Apple 提供的OpenCL platform，但是在其它系統(tǒng)上，可能會有不同廠商提供的多個不同的OpenCL platform，因此需要先取得platform 的數(shù)目：
cl_int err;
cl_uint num;
err = clGetPlatformIDs(0, 0, &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platforms\n";
return 0;
}
大部份的OpenCL API 會傳回錯誤值。如果傳回值是CL_SUCCESS 則表示執(zhí)行成功，否則會傳回某個錯誤值，表示失敗的原因。
接著，再取得platform 的ID，這在建立OpenCL context 時會用到：
std::vector<cl_platform_id> platforms(num);
err = clGetPlatformIDs(num, &platforms[0], &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platform ID\n";
return 0;
}
在OpenCL 中，類似這樣的模式很常出現(xiàn)：先呼叫第一次以取得數(shù)目，以便配置足夠的記憶體量。接著，再呼叫第二次，取得實際的資料。
接下來，要建立一個OpenCL context。如下：
cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };
cl_context context = clCreateContextFromType(prop, CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL);if(context == 0) {
std::cerr << "Can't create OpenCL context\n";return 0;
}clReleaseContext(context);
return 0; 在上面的程式中，clCreateContextFromType是一個OpenCL的API，它可以從指定的裝置類別中，建立一個OpenCL context。第一個參數(shù)是指定context的property。在OpenCL中，是透過一個property的陣列，以「property種類」及「property內(nèi)容」成對出現(xiàn)，并以0做為結束。例如，以上面的例子來說，要指定的property種類是CL_CONTEXT_PLATFORM，即要使用的platform ID，而property內(nèi)容則是由之前取得的platform ID中的第一個（即platforms[0]）。由于property的內(nèi)容可能是不同的資料型態(tài)，因此需要使用reinterpret_cast來進行強制轉型。
第二個參數(shù)可以指定要使用的裝置類別。目前可以使用的類別包括：
- CL_DEVICE_TYPE_CPU：使用CPU 裝置
- CL_DEVICE_TYPE_GPU：使用顯示晶片裝置
- CL_DEVICE_TYPE_ACCELERATOR：特定的OpenCL 加速裝置，例如CELL
- CL_DEVICE_TYPE_DEFAULT：系統(tǒng)預設的OpenCL 裝置
- CL_DEVICE_TYPE_ALL：所有系統(tǒng)中的OpenCL 裝置
這里使用的是CL_DEVICE_TYPE_DEFAULT，也就是指定使用預設的裝置。另外，在這里，直接使用了之前取得的OpenCL platform ID中的第一個ID（實際的程式中，可能會需要讓使用者可以指定要使用哪一個platform）。
如果建立OpenCL context失敗，會傳回0。因此，要進行檢查，并顯示錯誤訊息。如果建立成功的話，在使用完后，要記得將context釋放。這可以透過呼叫clReleaseContext來達成。
這個程式基本上已經(jīng)可以編譯執(zhí)行了，但是當然它并沒有真的做什么事情。
一個OpenCL context中可以包括一個或多個裝置，所以接下來的工作是要取得裝置的列表。要取得任何和OpenCL context相關的資料，可以使用clGetContextInfo函式。以下是取得裝置列表的方式：
size_t cb;
clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);
std::vector<cl_device_id> devices(cb / sizeof(cl_device_id));
clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0);
CL_CONTEXT_DEVICES表示要取得裝置的列表。和前面取得platform ID的情形相同，clGetContextInfo被呼叫了兩次：第一次是要取得需要存放裝置列表所需的記憶體空間大小（也就是傳入&cb），然后第二次呼叫才真正取得所有裝置的列表。接下來，可能會想要確定倒底找到的OpenCL裝置是什么。所以，可以透過OpenCL API取得裝置的名稱，并將它印出來。取得和裝置相關的資料，是使用clGetDeviceInfo函式，和前面的clGetContextInfo函式相當類似。以下是取得裝置名稱的方式：
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb);
std::string devname;
devname.resize(cb);
clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0);
std::cout << "Device: " << devname.c_str() << "\n";
到目前為止，完整的程式應該如下所示：// OpenCL tutorial 1
#include <iostream>
#include <string>
#include <vector>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
int main()
{
cl_int err;
cl_uint num;
err = clGetPlatformIDs(0, 0, &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platforms\n";
return 0;
}
std::vector<cl_platform_id> platforms(num);
err = clGetPlatformIDs(num, &platforms[0], &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platform ID\n";
return 0;
}
cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms[0]), 0 };
cl_context context = clCreateContextFromType( prop , CL_DEVICE_TYPE_DEFAULT, NULL, NULL, NULL); if(context == 0) { std::cerr << "Can't create OpenCL context\n"; return 0; } size_t cb; clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb); std::vector<cl_device_id> devices(cb / sizeof(cl_device_id)); clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, &devices[0], 0); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &cb); std::string devname; devname.resize(cb); clGetDeviceInfo(devices[0], CL_DEVICE_NAME, cb, &devname[0], 0); std::cout << "Device: " << devname.c_str() << "\n"; clReleaseContext(context); return 0;
}
執(zhí)行這個程式，如果建立OpenCL context 成功的話，應該會顯示出找到的OpenCL 裝置的名稱，例如
Device: GeForce GTX 285
建立Command Queue
大部份OpenCL 的操作，都要透過command queue。Command queue 可以接收對一個OpenCL 裝置的各種操作，并按照順序執(zhí)行（OpenCL 也容許把一個command queue 指定成不照順序執(zhí)行，即out-of-order execution，但是這里先不討論這個使用方式）。所以，下一步是建立一個command queue：
cl_command_queue queue = clCreateCommandQueue(context, devices[0], 0, 0);if(queue == 0) {
std::cerr << "Can't create command queue\n";clReleaseContext(context);
return 0;
} 和context 一樣，在程式結束前，要把command queue 釋放，即：clReleaseCommandQueue(queue);
上面的程式中，是把裝置列表中的第一個裝置（即devices[0]）建立command queue。如果想要同時使用多個OpenCL裝置，則每個裝置都要有自己的command queue。產(chǎn)生資料
由于這個程式的目的是要把一大堆數(shù)字進行相加，所以需要產(chǎn)生一些「測試資料」：
const int DATA_SIZE = 1048576;
std::vector<float> a(DATA_SIZE), b(DATA_SIZE), res(DATA_SIZE);
for(int i = 0; i < DATA_SIZE; i++) {
a[i] = std::rand();
b[i] = std::rand();
}配置記憶體并復制資料
要使用OpenCL 裝置進行運??算時，通常會需要在OpenCL 裝置上配置記憶體，并把資料從主記憶體中復制到裝置上。有些OpenCL 裝置可以直接從主記憶體存取資料，但是速度通常會比較慢，因為OpenCL 裝置（例如顯示卡）通常會有專用的高速記憶體。以下的程式配置三塊記憶體：
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);
cl_mem cl_res = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * DATA_SIZE, NULL, NULL);
if(cl_a == 0 || cl_b == 0 || cl_res == 0) {
std::cerr << "Can't create OpenCL buffer\n";
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}clCreateBuffer函式可以用來配置記憶體。它的第二個參數(shù)可以指定記憶體的使用方式，包括：- CL_MEM_READ_ONLY：表示OpenCL kernel 只會對這塊記憶體進行讀取的動作
- CL_MEM_WRITE_ONLY：表示OpenCL kernel 只會對這塊記憶體進行寫入的動作
- CL_MEM_READ_WRITE：表示OpenCL kernel 會對這塊記憶體進行讀取和寫入的動作
- CL_MEM_USE_HOST_PTR：表示希望OpenCL 裝置直接使用指定的主記憶體位址。要注意的是，如果OpenCL 裝置無法直接存取主記憶體，它可能會將指定的主記憶體位址的資料復制到OpenCL 裝置上。
- CL_MEM_ALLOC_HOST_PTR：表示希望配置的記憶體是在主記憶體中，而不是在OpenCL 裝置上。不能和CL_MEM_USE_HOST_PTR 同時使用。
- CL_MEM_COPY_HOST_PTR：將指定的主記憶體位址的資料，復制到配置好的記憶體中。不能和CL_MEM_USE_HOST_PTR 同時使用。
第三個參數(shù)是指定要配置的記憶體大小，以bytes為單位。在上面的程式中，指定的大小是sizeof(cl_float) * DATA_SIZE。
第四個參數(shù)是指定主記憶體的位置。因為對cl_a和cl_b來說，在第二個參數(shù)中，指定了CL_MEM_COPY_HOST_PTR，因此要指定想要復制的資料的位址。cl_res則不需要指定。
第五個參數(shù)是指定錯誤碼的傳回位址。在這里并沒有使用到。
如果clCreateBuffer因為某些原因無法配置記憶體（例如OpenCL裝置上的記憶體不夠），則會傳回0。要釋放配置的記憶體，可以使用clReleaseMemObject函式。
編譯OpenCL kernel 程式
現(xiàn)在執(zhí)行OpenCL kernel 的準備工作已經(jīng)大致完成了。所以，現(xiàn)在剩下的工作，就是把OpenCL kernel 程式編釋并執(zhí)行。首先，先把前面提過的OpenCL kernel 程式，存放在一個文字檔中，命名為shader.cl：
__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];
}要編譯這個kernel程式，首先要把檔案內(nèi)容讀進來，再使用clCreateProgramWithSource這個函式，然后再使用clBuildProgram編譯。如下所示：
cl_program load_program(cl_context context, const char* filename)
{
std::ifstream in(filename, std::ios_bas??e::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];
cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);
if(program == 0) { return 0;}
if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) { return 0;}
return program;} 上面的程式，就是直接將檔案讀到記憶體中，再呼叫clCreateProgramWithSource建立一個program object。建立成功后，再呼叫clBuildProgram函式編譯程式。clBuildProgram函式可以指定很多參數(shù)，不過在這里暫時沒有使用到。
有了這個函式，在main 函式中，直接呼叫：
cl_program program = load_program(context, "shader.cl");
if(program == 0) {
std::cerr << "Can't load or build program\n";
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
同樣的，在程式結束前，要記得將program object 釋放：clReleaseProgram(program);
一個OpenCL kernel 程式里面可以有很多個函式。因此，還要取得程式中函式的進入點：
cl_kernel adder = clCreateKernel(program, "adder", 0);
if(adder == 0) {
std::cerr << "Can't load kernel\n";
clReleaseProgram(program);
clReleaseMemObject(cl_a);
clReleaseMemObject(cl_b);
clReleaseMemObject(cl_res);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}和program object 一樣，取得的kernel object 也需要在程式結束前釋放：clReleaseKernel(adder);
執(zhí)行OpenCL kernel
弄了這么多，總算可以執(zhí)行OpenCL kernel程式了。要執(zhí)行kernel程式，只需要先設定好函式的參數(shù)。adder函式有三個參數(shù)要設定：
clSetKernelArg(adder, 0, sizeof(cl_mem), &cl_a);
clSetKernelArg(adder, 1, sizeof(cl_mem), &cl_b);
clSetKernelArg(adder, 2, sizeof(cl_mem), &cl_res);
設定參數(shù)是使用clSetKernelArg函式。它的參數(shù)很簡單：第一個參數(shù)是要設定的kernel object，第二個是參數(shù)的編號（從0開始），第三個參數(shù)是要設定的參數(shù)的大小，第四個參數(shù)則是實際上要設定的參數(shù)內(nèi)部。以這里的adder函式來說，三個參數(shù)都是指向memory object的指標。設定好參數(shù)后，就可以開始執(zhí)行了。如下：
size_t work_size = DATA_SIZE;
err = clEnqueueNDRangeKernel(queue, adder, 1, 0, &work_size, 0, 0, 0, 0);
clEnqueueNDRangeKernel會把執(zhí)行一個kernel的動作加到command queue里面。第三個參數(shù)（1）是指定work item數(shù)目的維度，在這里就是一維。第五個參數(shù)是指定work item的總數(shù)目，也就是DATA_SIZE。后面的參數(shù)現(xiàn)在暫時先不用管。如果成功加入的話，會傳回CL_SUCCESS。否則會傳回錯誤值。在執(zhí)行kernel 被加到command queue 之后，就可能會開始執(zhí)行（如果command queue 現(xiàn)在沒有別的工作的話）。但是clEnqueueNDRangeKernel 是非同步的，也就是說，它并不會等待OpenCL 裝置執(zhí)行完畢才傳回。這樣可以讓CPU 在OpenCL 裝置在進行運算的同時，進行其它的動作。
由于執(zhí)行的結果是在OpenCL 裝置的記憶體中，所以要取得結果，需要把它的內(nèi)容復制到CPU 能存取的主記憶體中。這可以透過下面的程式完成：
if(err == CL_SUCCESS) {
err = clEnqueueReadBuffer(queue, cl_res, CL_TRUE, 0, sizeof(float) * DATA_SIZE, &res[0], 0, 0, 0);
}clEnqueueReadBuffer函式會把「將記憶體資料從OpenCL裝置復制到主記憶體」的動作加到command queue中。第三個參數(shù)表示是否要等待復制的動作完成才傳回，CL_TRUE表示要等待。第五個參數(shù)是要復制的資料大小，第六個參數(shù)則是目標的位址。由于這里指定要等待復制動作完成，所以當函式傳回時，資料已經(jīng)完全復制完成了。最后是進行驗證，確定資料正確：
if(err == CL_SUCCESS) {
bool correct = true;
for(int i = 0; i < DATA_SIZE; i++) {
if(a[i] + b[i] != res[i]) {
correct = false;
break;
}}if(correct) { std::cout << "Data is correct\n"; }
else {
std::cout << "Data is incorrect\n";
}}
else { std::cerr << "Can't run kernel or read back data\n";} 到這里，整個程式就算是完成了。編譯后執(zhí)行，如果順利的話，應該會印出
Data is correct
的訊息。以下是整個程式的全貌：
// OpenCL tutorial 1
#include <iostream>
#include <fstream>
#include <string>
#include <vector>
#include <cstdlib>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
cl_program load_program(cl_context context, const char* filename)
{
std::ifstream in(filename, std::ios_bas??e::binary); if(!in.good()) { return 0; } // get file length in.seekg(0, std::ios_bas??e::end); size_t len??gth = in.tellg(); in.seekg(0, std::ios_bas??e::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]; cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0); if(program == 0) { return 0; } if(clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS) { return 0; } return program;}
int main()
{
cl_int err;
cl_uint num;
err = clGetPlatformIDs(0, 0, &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platforms\n";
return 0;
}
std::vector<cl_platform_id> platforms(num);
err = clGetPlatformIDs(num, &platforms[0], &num);
if(err != CL_SUCCESS) {
std::cerr << "Unable to get platform ID\n";
return 0;
}
}
在附件中可以下載包括Xcode project 以及Visual Studio 2008 project 檔的原始碼。
Attachments ( 1 )
- cltut_1.zip - on Feb 3, 2010 8:54 AM by Chen Ping-Che (version 2 / earlier versions ) 7k Download
登錄 最近的站點活動 服務條款 舉報不良信息 打印頁面 由Google協(xié)作平臺強力驅動
轉載于:https://www.cnblogs.com/wangshide/articles/2218354.html
總結
以上是生活随笔為你收集整理的[转]OpenCL 教学（一）的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： 自定义EL函数解决JSTL标签不足之处—
- 下一篇： 最近忙，三个字