背景

最近在學OpenCL，用的書是《OpenCL異構並行編程實戰》。怎麼說呢，感覺這本書比較迷，講的很亂，跟著看完的話可能學不到什麼。好在之前學過CUDA，勉強能夠理解並行計算的思路。因此這裡寫了一下用顯卡進行RGB2Gray的程序，也算是自己能夠使用OpenCL進行一些簡單的顯卡計算吧。

平台：Ubuntu18.04 + CUDA10.0 + MX150 + OpenCL 1.2

目標：從磁碟讀取一張圖片，送入OpenCL設備進行RGB轉灰度圖，再拷貝回來顯示

使用OpenCL進行計算的流程

無論是OpenCL，還是CUDA，當利用顯卡計算時，都需要經歷如下步驟：

1. 設備初始化 2. 準備主機端數據(分配主機端內存+獲取數據) 3. 分配設備端內存 4. 將主機端數據拷貝到設備端 5. 設備啟動內核函數，進行運算，將結果寫到設備端內存 6. 將設備端結果拷貝回主機端 7. 讀取主機端內存，進行後續處理 8. 釋放資源

對於OpenCL，具體的步驟是：

1. 設備初始化(獲取平台和設備id，創建上下文和命令隊列) 2. 編寫並編譯內核 3. 準備主機端數據並傳入設備(準備主機端數據，創建設備端緩衝對象，傳入數據) 4. 啟動內核函數(傳遞參數，啟動內核) 5. 將結果拷貝回主機端 6. 後續處理 7. 釋放資源

設備初始化

獲取平台id

cl_int error;
cl_platform_id platform;

error = clGetPlatformIDs(1, &platform, NULL); //獲取平台id

這裡clGetPlatformIDs的函數原型為：

clGetPlatformIDs(cl_uint /* num_entries */,
cl_platform_id * /* platforms */,
cl_uint * /* num_platforms */)

本例是只讀取一個平台。實際上可能有多個平台。因此實際上更為通用的使用方法為：

cl_uint numOfPlatforms;
cl_int error;

error = clGetPlatformIDs(0, NULL, &numOfPlatforms); //獲取平台數量
if(error != CL_SUCCESS)
{
perror("Cannot get platform ids");
exit(1);
}
cl_platform_id *platforms = (cl_platform_id*)alloca(sizeof(cl_platform_id)*numOfPlatforms);
error = clGetPlatformIDs(numOfPlatforms, platforms, NULL); //獲取cl_platform_id實體

這裡我們看到clGetPlatformIDs()函數被調用了兩次，第一次是獲取平台的數量numOfPlatforms，第二次是獲取cl_platform_id的實體。OpenCL里有很多類似的API。

本例是為了偷個懶，只使用第一個平台。

獲取設備id

獲取完平台id後，要獲取設備id。這裡一樣偷懶，只使用平台上第一個設備。實際上一個平台上可能有多個OpenCL設備(CPU、GPU)。

cl_device_id device;

error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL); //獲取設備id

函數原型為：

clGetDeviceIDs(cl_platform_id /* platform */,
cl_device_type /* device_type */,
cl_uint /* num_entries */,
cl_device_id * /* devices */,
cl_uint * /* num_devices */)

創建設備上下文

OpenCL使用上下文來管理設備，因此不論進行什麼操作，都需要先創建上下文：

context = clCreateContext(NULL, 1, &device, NULL, NULL, &error); //創建上下文

創建命令隊列

命令隊列和CUDA中的流的概念相對應。主機程序在設備上創建命令隊列，並向命令隊列中壓入操作(數據傳輸、內核執行)。 一個設備上可以創建多個命令隊列。

與CUDA略有區別的是，OpenCL除了順序執行命令隊列中的任務之外，還可以不按順序執行(創建隊列時傳入CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE標誌位)。

創建命令隊列時，需要傳入設備id和上下文

cQ = clCreateCommandQueue(context, device, NULL, &error);

函數原型：

clCreateCommandQueue(cl_context /* context */,
cl_device_id /* device */,
cl_command_queue_properties /* properties */,
cl_int * /* errcode_ret */)

編譯內核

與CUDA不同的是，OpenCL需要在程序中顯式編譯內核。

內核函數

OpenCL工程一般將內核放在單獨的*.cl文件中。本例的內核為：

/*rgb2gray.cl*/
__kernel void kernel_rgb2gray(__global unsigned char * rgbImage,
__global unsigned char * grayImage,
__global unsigned * const p_height,
__global unsigned * const p_width)
{
int x = get_global_id(0);
int y = get_global_id(1);
int height = *p_height;
int width = *p_width;
if(x < width && y < height)
{
int index = y * width + x;
grayImage[index] = 0.299f*rgbImage[index*3] +
0.587f*rgbImage[index*3+1] +
0.114f*rgbImage[index*3+2];
}
}

程序很簡單，每個工作項(對應於CUDA中的線程)讀取一個像素，計算灰度值，寫回該像素。

由於圖像尺寸不定，因此工作項不可能剛好和圖像尺寸相對匹配，因此需要if(x < width && y < height)來分支。height和width則是由主機程序啟動內核時傳入的。

編譯內核

OpenCL需要顯示編譯內核，其流程為：

讀取內核文件->創建program對象->編譯程序->創建內核

讀取內核文件的代碼如下：

const char *file_names[] = {"rgb2gray.cl"}; //待編譯的內核文件
 const int NUMBER_OF_FILES = 1;
char* buffer[NUMBER_OF_FILES];
size_t sizes[NUMBER_OF_FILES];

loadProgramSource(file_names, NUMBER_OF_FILES, buffer, sizes); //讀取內核文件文本
 program = clCreateProgramWithSource(context, NUMBER_OF_FILES, (const char**)buffer, sizes, &error); //創建program對象
 error = clBuildProgram(program, 1, &device, NULL, NULL, NULL); //編譯程序

if(error != CL_SUCCESS) {
// If theres an error whilst building the program, dump the log
 size_t log_size;
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
char *program_log = (char*) malloc(log_size+1);
program_log[log_size] =

相关文章