OpenCL 介紹

前言

在以往只有 fixed function pipeline 的時代, GPU 單純做為 3D 繪圖加速器, 應用程式透過諸如 OpenGL, OpenGL ES, Direct3D 等等的 3D API, 使用硬體提供的繪圖加速服務, 在 programmable pipeline 出現之後, GPU的應用進入了另一個領域 : computing. 這個階段出現了一個新的名詞 : GPGPU; General Purpose Computing on Graphics Processing Unit. GPU 內的 shader, texturing pipeline, arithmetic pipeline以及 Multi-thread 架構, 開始被應用在計算領域. 在這個階段, 計算應用必須使用標準的 3D API來完成, 這對計算應用在實作上造成不便, 例如, 要實做一個簡單的浮點數資料相加計算程式, 需要涉及到像是 vertex data, texture 等 和計算無關的資料, 概念和流程.
為了讓 GPU 更方便使用在計算領域, 專門為 computing 而制定的 API 出現了,
這讓 GPU 在 computing 的應用進入了第二個階段, 這個階段的概念是 Heterogeneous Computing; 異質性計算, 要能使用系統上不同性質的運算單元, 舉例來說, 在 CPU與 GPU 的協同運算中, 由CPU 負責程式的流程控制, 大量的資料則交由 GPU 以平行化的方式進行計算.
常見的應用包括:
image processing
computer vision
speech processing

可以參考筆者另一篇文章 OpenCL 毛玻璃效果

目前業界主要的 compute API 有以下四種:
(1) OpenCL ; Apple 首先提出, 並聯合多家廠商合組 working group, 目前由 Khronos 維護. 是本文章要介紹的API.
(2)C++AMP, DirectCompute ; 由 MicroSoft 主導制定.
(3)CUDA ; nVidia 自家的規格.
(4)Renderscript ; Google 制定, Android 3.0 開始出現.

其實使用這些 API 的目的不外乎就是要藉 GPU 的平行運算能力, 縮短計算任務所需的時間.
因此以平行化運算的方式思考, 應該是使用 GPU 加速運算的重點.

Laplace filter 影像處理

傳統的循序式計算如下 :
for(int x=0 ; x<width;++x)
  for(int y=0 ; y<height;++y){
     // filtering pixel(x,y)
     dst_img[x][y] = LF(src_img[x][y]);
  }

以圖二和圖三間單說明 循序式處理和平行處理的差異.

圖二 循序式處理

圖三 平行式處裡

OpenCL
OpenCL 運作的基本概念, 可從下面四個面向了解:
 (1) Platform model.
 (2) Execution model.
 (3) Memory model.
 (4) Programming model

OpenCL Platform model 

一個 OpenCL 應用程式包括兩個部份, 一部分運行在 GPU 上的平行運算程式, 稱為 kernel, 

是由 OpenCL C languange 語法撰寫. 另一部分運行在 Host(CPU) 的程式, 負責 流程控制, 例如:

    (1) compile OpenCL kernel 

    (2) 建立 memory buffer 

    (3) 設定 command queue  

    (4) 設定 ND range 

OpenCL使用 Platform model 描述 GPU 運算資源, Platform model 包括幾個名詞

   (1) kernel:這在前面解釋過, 是在 GPU 上運行的平行程式.

   (2) work-item : 這是運行 kernel 程式的最小單位, 通常對應 GPU 的 thread.

   (3) work group: 一群 work-item 組成一個 work group, 不同系列的 GPU , 一個 work group 所包

        括 的最大數量 work-item 會有不同. 這跟GPU 的memory resource 以及 kernel code 使用變
        量的數目有關係, 基本上, kernel code 使用的變量越多, work group 包括的 work item 數量
        會越少, GPU 的 shader core 是以 work group 為單位進行 schedule.  work group 對應到
         OpenCL 的 compute element.

   (4) compute unit : 一個 compute unit 包括多個 compute element.

   (5) OpenCL device : 一個 OpenCL device 包括多個 compute unit.一個系統中可能包括多個 

         OpenCL device.

 下面用圖四 表示這些 OpenCL 名詞和 GPU 架構之間的對應關係.  不同的架構可能有不同的對應關係.

圖四OpenCL Platform model 和 GPU 架構對應關係

OpenCL Execution model

Execution model 決定 kernel 會如何被運行, 也就是 work-item 和 待處理資料的對應關係, 以影像處理當做例子, 一個 work-item 可以處理一個 pixel 或是多個 pixel 的運算. 若是處理一個 pixel, 則 workitem 對一個 pixel 執行 kernel 所定義的運算. 

OpenCL 使用 ND-Range 來定義資料的切割方式, 也就是資料要切割成多少個 work-item 處理. OpenCL 透過 global_work_size 定義work-item 的數量, OpenCL 的 global_work_size 可以是一維, 二維 或是三維.

圖五

如圖五上方所示.一個一維資料陣列有 10 筆資料, 若是設定 global_work_size[1] = 10.

則總共有10 筆 workitem. 圖五下方一個8x8 的二維資料陣列, 若設定global_work_size[2]={8,8}, 則總共有 64 個work-item. 若是設定 global_work_size[2]={8/2, 8}, 則總共會有 32 個 work-item, 每個 work-item 處理兩筆資料.

另外, host program 也可以透過設定 local_work_size 來決定 work group 的 size, 以圖五的二維資料為例, 若設定 local_work_size[2] = {4,4}, 則每個 work group 由 16 個 work-item 組成, 且總共有 (8x8) / (4x4) = 4 個 work group. 如圖六所示.

圖六 4個 work group

標示為橘色的 work-item 在 ND-Range 中的位置為如下:

get_global_id(0) = 3
get_global_id(1) = 6

get_local_id(0) = 3
get_local_id(1) = 2

get_local_size(0) = 4
get_local_size(1) = 4

get_group_id(0) = 0
get_group_id(1) = 1

在同一個 work group 中的 所有 work-item 可以進行同步且共用 shared memory(請參考 OpenCL memory model).
GPU 的架構是所謂 MTTOP ;Massively-Thread Throughput-Oriented Processor,
每個 work-item 對應到 GPU shader core 中的一條 thread. thread scheduler 以 SIMT(Single Instruction Multiple Thread)的方式讓每個 shader core 處理大量的 thread . 其概念是準備儘量多的 thread 讓 GPU 的 shader core 進行處理, 當執行中的一組 thread 因為 I/O (例如:memory access) 得等待(通常需要數個 clock cycle)而停止往下執行時, 可以切換到其他組正在準備執行的 thread, 如此隱藏住I/O 造成的延遲(latency hiding), 讓 shader core 內部的 ALU/FPU 等運算資源始終保持較高的利用率以提高 throughput.
以 nVidia 目前的 GPU 架構為例, 其 thread scheduler 是以 warp (32 個 threads)為單位, warp 中的 threads 共用同一個 program counter, 若 kernel的 程式中有 if-then-else, switch, for loop 等等 的 branch 發生時, 需要循序地將不同 branch 的 thread 執行完畢, 此情況稱為 thread divergence, 有些廠商的 GPU 則是每個 thread 有自己的 program counter 來更有效地處理 thread divergence.

圖七 Execution branch

圖八 Thread Divergence

OpenCL Memory Model

OpenCL 規範了四種記憶體, 如圖九所示.

圖九 OpenCL Memory Model

(1) private memory : 只供 work-item 自己讀寫.
(2) local memory : 在同一個 work group 中的所有 work item 可共用.
(3) constant memory : 唯讀的記憶體, 全部的 work-item 皆可使用.
(4) global memory : 全部 work-item 皆可讀寫.


接下來用 Laplacian filter 做為本文的程式範例. 使用 OpenCV 實作host版本的 Laplacian filter 和使用 OpenCL 實作的 GPU 版本比較執行 filtering 所需時間.

我們先來看 host 這一邊的程式代碼.
代碼的大致流程如下:
(1.1)OpenCV 讀入圖檔
(1.2)使用 OpenCV 的 filter2D, 對圖檔進行 laplacian filtering.
(1.3)紀錄步驟(1.2)運行所花費的時間.
(2.1)OpenCL 初始化 : 包括初始化 platform, context, device, command queue 等等.
(2.2)準備資料 : 在 host 端把影像資料準備好, 本文使用 OpenCV 載入圖檔以及顯示 GPU
       處理後的影像.
(2.3)將host 端的影像資料 複製到 GPU.
(2.4)GPU 執行 kernel program.
(2.5)將GPU 執行的結果讀回 host, 使用 OpenCV 顯示.

 步驟(1.1)~(1.3)的程式代碼如下:
 Mat gray_Img0, gray_Img1;
Mat CPU_img;
Mat Img1 = cv::imread("0000.jpg");
if( Img1.empty())
return 0;
Mat Img0;
cvtColor(Img1, Img0, CV_RGB2RGBA);
//Cvt_2_Gray(Img0, gray_Img0);
int ch = Img0.channels();

imshow( "original image\n ", Img0 );
printf("img_w=%d, img_h=%d\n", Img0.cols, Img0.rows);

int kernel_size = 3;
Mat filter_kernel(kernel_size, kernel_size, CV_8S);
Point anchor = Point(-1,-1);
double delta;
int ddepth;
filter_kernel.data[0] = -1;  filter_kernel.data[1] = -1;  filter_kernel.data[2] = -1;
filter_kernel.data[3] = -1;  filter_kernel.data[4] = 9   filter_kernel.data[5] = -1;
filter_kernel.data[6] = -1;  filter_kernel.data[7] = -1;  filter_kernel.data[8] = -1;
QueryPerformanceCounter(&t1);
filter2D( Img0, CPU_img, -1, filter_kernel, Point(-1,-1) );
QueryPerformanceCounter(&t2);
CPU_elapse = (double)(t2.QuadPart-t1.QuadPart)/(double)(ts.QuadPart);
printf("CPU Time consume: %lf\n", CPU_elapse );

imshow( "CPU image", CPU_img );

接下來看OpenCL 的初始化, 首先是取得系統上的 OpenCL platform, 筆者系統上只有一個 Nvidia OpenCL platform, 因此 clGetPlatform 的第一個參數直接指定為 1(第一個 platform).
cl_int err;
cl_uint num = 0;
cl_platform_id platforms = 0;

if (!checkSuccess(clGetPlatformIDs(1, &platforms, &num)))
{
      cerr << "Retrieving OpenCL platforms failed. " << __FILE__ << ":"<< __LINE__ << endl;
      return false;
}

得到系統的 platform 後, 可以建立一個 OpenCL context.

cl_context_properties prop[] = { CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties>(platforms), 0 };
GPU_context = clCreateContextFromType(prop, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
if(GPU_context == 0) {
cerr << "Can't create OpenCL context\n";
return 0;
}
clCreateContextFromType 第二個參數表示我們要為一個 GPU device 建立 context, 第一個參數是 context 的 property, 型別是 cl_context_properties, OpenCL 1.1 目前只有CL_CONTEXT_PLATFORM 一個 property, cl_context_properties 的內容是一個 property 對應一個 value.

接著我們調用 clGetContextInfo 取得 device ID, 因為系統中可能有不只一個 device, 所以共調用 clGetContextInfo 兩次, 第一次取得 device 的數量, 第二次才取得所有的 device id.
size_t deviceBUfSize;
clGetContextInfo(GPU_context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBUfSize);
vector<cl_device_id> devices(deviceBUfSize / sizeof(cl_device_id));
clGetContextInfo(GPU_context, CL_CONTEXT_DEVICES, deviceBUfSize, &devices[0], 0);

有了 device id 之後, 我們可以查詢一些 device 的資訊, 例如 device 有幾個 compute unit. 一個 work group 最大支持的 work-item 數量.
// get device compute unit.
cl_uint compute_unit;
clGetDeviceInfo(devices[0], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &compute_unit, 0);
// get max. work group size.
cl_uint max_work_group_size;
clGetDeviceInfo(devices[0], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(cl_uint), &max_work_group_size, 0);

在本文中, 我們要詢問的重點是device 是否支持 image.
cl_bool img_support;
clGetDeviceInfo(devices[0], CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &img_support, 0);
if( img_support == CL_TRUE )
printf("device support image\n");

接下來, 要建立 command queue, command queue 是 host 下達命令到 device 的管道, 例如
要 device 執行 kernel 的命令或是 將資料從 device 讀回 host 的命令 或是 從 host 寫資料到 device 的命令. 如圖十所示:

圖十Host 透過 command queue 將 command 送往 Device

commandQueue = clCreateCommandQueue(GPU_context, devices[0], 0, 0);
if(commandQueue == 0) {
std::cerr << "Can't create command queue\n";
clReleaseContext(GPU_context);
return 0;
}

影像資料需要從Host 送到 Device, 經過 Device 處理後, 再 copy 回 Host, 因此要建立兩個 image object, 跟別供 Device read以及 write. cl_image_format 用來設定 image 的 color format, 在此例中, 我們設定 color format 為 RGBA8888:
cl_image_format input_img_fmt;
input_img_fmt.image_channel_order = CL_RGBA;
input_img_fmt.image_channel_data_type = CL_UNSIGNED_INT8;

cl_image_format output_img_fmt;
output_img_fmt.image_channel_order = CL_RGBA;
output_img_fmt.image_channel_data_type = CL_UNSIGNED_INT8;

// create image object
imgObject[0] = clCreateImage2D(GPU_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
                                                &input_img_fmt, Img0.cols, Img0.rows, 0, Img0.data, NULL);
imgObject[1] = clCreateImage2D(GPU_context, CL_MEM_WRITE_ONLY , 
                                                &output_img_fmt, Img0.cols, Img0.rows, 0, NULL, NULL);

第一個 clCreateImage2D 建立的 image eobject 要給 device 讀取, 所以第二個參數使用 CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, host 會將 Img0.data 的內容 copy 到此 image object. 第二個 image object 則是 device 將運算結果寫出用的.

接著建立 program, 將 Device 端的程式代碼從檔案讀入並且  compile.
char shader_name[] = "filter_shader_img.cl";
printf("%s\n", shader_name);
program = load_program( GPU_context, shader_name, devices[0] );
if( program == 0 )
{
cleanUpOpenCL(GPU_context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
return 0;
}
load_program 函式的內容會在後面的程式中列出.

建立 kernel object, 第一個參數是剛剛建立的 program object, 第二個參數是 kernel 的函式名稱.
kernel = clCreateKernel( program, "laplacian", NULL );
if( kernel == 0 )
{
cleanUpOpenCL(GPU_context, commandQueue, program, kernel, memoryObjects, numberOfMemoryObjects);
return 0;

}
接著 設定要傳入 kernel 函式的參數, 有兩個參數, 是前面建立好的 imgObject[0] 以及 imgObject[1]. 分別對應 laplacian kernel 函式的參數 0 以及參數 1:

clSetKernelArg(kernel, 0, sizeof(cl_mem), &imgObject[0]);  // source image
clSetKernelArg(kernel, 1, sizeof(cl_mem), &imgObject[1]);  // output image

設定 global work size, 在這裡, 我們將 global work size 設定成 image的寬以及高, 並且將 執行 kernel 的命令發到 command queue:
const int workDimensions = 2;
size_t globalWorkSize[workDimensions] = {Img0.cols, Img0.rows}; 

cl_event evt, r_evt;
err = clEnqueueNDRangeKernel(commandQueue, kernel, workDimensions, NULL, globalWorkSize, NULL, 0, NULL, &evt);

size_t origin[3] = {0,0,0};
size_t region[3] = {Img0.cols, Img0.rows, 1};
err = clEnqueueReadImage(commandQueue, imgObject[1], CL_FALSE, origin, region, 0, 0, &res[0], 1, &evt, &r_evt);

clWaitForEvents( 1, &r_evt );
Mat ocl_img(Img0.rows, Img0.cols, CV_8UC4, &res[0]);

imshow("GPU image", ocl_img );

clEnqueueNDRangeKernel 將執行 kernel 的命令發到 command queue, 在 OpenCL  specification 中沒有明確規範這個函式調用是 blocking 或是 non-blocking,  要視各家的實做而定.  OpenCL 提供了豐富的 event 機制 可做 command 之間的同步 或是 host 和  device 之間的同步, 此外也可以做為 profiling 使用. 在 clEnqueueNDRangeKernel 最後一個參數給定一個 OpenCL event : evt 當 device 運算完所有資料後 會更新 evt 的狀態.
clEnqueueReadImage 的作用是將 device 運算完成的結果讀回 host. 在這裡, 我們需要確定 device 已經處理完畢了, 才能將處理的結果讀回, 因此 clEnqueueNDRangeKernel 使用 event : evt, 來通知 command queue 中的clEnqueueReadImage 命令是否可以被執行.  clEnqueueReadImage 的倒數第三個參數是 Event list 中event 的數目,倒數第二個參數是一個 event list,  表示該命令被執行前需要等待的 Event 哪些, 在這個例子, 要等待一個由 clEnqueueNDRangeKernel產生的 event : evt. 此外, 這邊將 clEnqueueReadImage 命令 第三個參數為 CL_FALSE, 代表將此函式的調用設定為 non-blocking, 這是比較有效率的使用方式, 如此在資料讀取的過程中, host 可以繼續執行其它的工作, 因為是 non-blocking 調用方式, 所以 clEnqueueReadImage 最後一個參數需要給定一個 event : r_evt. 做為同步使用.
host 這邊要調用 openCV 的 imshow 將結果顯示出來, 因此在此之前要確定資料已經從 device 讀回, 實做的方式是 host 這邊用 clWaitForEvent 等待.r_evt.
看完 host 端得程式代碼, 我們來看 device 端的程式代碼.

01 __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
02                                CLK_ADDRESS_CLAMP |
03                                CLK_FILTER_NEAREST;
04 
05 __kernel void laplacian( read_only  image2d_t src_image,
06                          write_only image2d_t dst_image)
07 {
08     int x = get_global_id(0);
09     int y = get_global_id(1);
10 
11     //0, 1, 2        -1, -1, -1
12     //3, 4, 5   =>   -1,  9, -1
13     //6, 7, 8        -1, -1, -1
14     int2 coord_0 = (int2)(x-1, y-1);
15     int2 coord_1 = (int2)(x,   y-1);
16     int2 coord_2 = (int2)(x+1, y-1);
17     int2 coord_3 = (int2)(x-1, y);
18     int2 coord_4 = (int2)(x,   y);
19     int2 coord_5 = (int2)(x+1, y);
20     int2 coord_6 = (int2)(x-1, y+1);
21     int2 coord_7 = (int2)(x,   y+1);
22     int2 coord_8 = (int2)(x+1, y+1);
23 
24     uint4 up_0 = read_imageui(src_image,sampler, coord_0);
25     uint4 up_1 = read_imageui(src_image,sampler, coord_1);
26     uint4 up_2 = read_imageui(src_image,sampler, coord_2);
27     uint4 up_3 = read_imageui(src_image,sampler, coord_3);
28     uint4 up_4 = read_imageui(src_image,sampler, coord_4);
29     uint4 up_5 = read_imageui(src_image,sampler, coord_5);
30     uint4 up_6 = read_imageui(src_image,sampler, coord_6);
31     uint4 up_7 = read_imageui(src_image,sampler, coord_7);
32     uint4 up_8 = read_imageui(src_image,sampler, coord_8);
33 
34     int4  sp_0 = convert_int4(up_0);
35     int4  sp_1 = convert_int4(up_1);
36     int4  sp_2 = convert_int4(up_2);
37     int4  sp_3 = convert_int4(up_3);
38     int4  sp_4 = convert_int4(up_4);
39     int4  sp_5 = convert_int4(up_5);
40     int4  sp_6 = convert_int4(up_6);
41     int4  sp_7 = convert_int4(up_7);
42     int4  sp_8 = convert_int4(up_8);
43 
44     int4 p_convo = (int4)0 - sp_0 -       sp_1     - sp_2
45                    - sp_3 + (int4)9*sp_4   - sp_5
46                    - sp_6 -       sp_7     - sp_8;
47     uint4 res = convert_uint4( clamp( p_convo, (int4)0, (int4)255 ) );
48     write_imageui(dst_image, coord_4, res);
49 }
50 

第 01~03 行,是建立 sampler, 在 OpenCL 中要使用 image, 除了前面提到的 建立 image object 之外, 另外就是 sampler, image object 代表影像資料的儲存方式, sampler 則代表 kernel 對影像資料的讀取方式.  sampler 需要定義三種特性, 分別是
(1) 座標是否 normalize : 座標數值是否正規化(0.0~1.0)
(2) Addressing mode : 若是座標超過 image 邊界時, 調整座標的方式, 目前有
      CLK_ADDRESS_CLAMP, CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_REPEAT,
      CLK_ADDRESS_MIRRORED_REPEAT, 以及 CLK_ADDRESS_NONE.

(3) Filtering mode : 座標調整後, 根據座標 對image 上的像素抓取時可以由單一像素或是多個像素決定最後點的顏色. 目前可用的選擇有 CLK_FILTER_NEAREST 或是
      CLK_FILTER_LINEAR.
   
有使用過一些 3D API ,例如 OpenGL  的朋友們對這些名詞應該不陌生. 可以參考 OpenGL 相關書籍, 筆者在這邊就不綴述.
OpenCL 可以將 sampler 定義在 host program 中, 也可以直接定義在 kernel program中, 筆者選擇後者, 因為這樣使用方式比較簡單.
08~09 行 每個 thread 使用 get_global_id 獲得存取 image 的座標.
14~22 行 因為 filter 的 mask 是 3X3, 共需要九組座標.
24~32 行, 調用 read_imageui 從 image 讀取九個點的顏色.
34~42 行, 從 unsigned 轉換成 signed.
44~46 行, 做 convolution.
47~48 行, 將顏色值限制在 0~255, 並且轉換回 unsigned 後調用 write_imageui 寫到 dest_image.

最後列出完整的 host program:

01 
02 #include "stdio.h"
03 #include <CL/cl.h>
04 #include <iostream>
05 #include <fstream>
06 
07 #include <vector>
08 #include <windows.h>
09 
10 #include <opencv2/features2d/features2d.hpp>
11 #include <opencv2/nonfree/features2d.hpp>
12 #include "opencv2/opencv.hpp"
13 
14 using namespace std;
15 using namespace cv;
16 #define XPRINTF(fmt,...) printf(fmt, __VA_ARGS__)
17 
18 inline bool
19 checkSuccess(cl_int errorNumber)
20 {
21     if (errorNumber != CL_SUCCESS)
22     {
23         //cerr << "OpenCL error: " << errorNumberToString(errorNumber) << endl;
24         return false;
25     }
26     return true;
27 }
28 
29 
30 bool
31 cleanUpOpenCL(cl_context context, cl_command_queue commandQueue,
32               cl_program program, cl_kernel kernel,
33               cl_mem* memoryObjects,
34               int numberOfMemoryObjects)
35 {
36     bool returnValue = true;
37     if (context != 0)
38     {
39         if (!checkSuccess(clReleaseContext(context)))
40         {
41             cerr << "Releasing the OpenCL context failed. " << __FILE__ <<
42                  ":"<< __LINE__ << endl;
43             returnValue = false;
44         }
45     }
46 
47     if (commandQueue != 0)
48     {
49         if (!checkSuccess(clReleaseCommandQueue(commandQueue)))
50         {
51             cerr << "Releasing the OpenCL command queue failed. " << __FILE__ <<
52                  ":"<< __LINE__ << endl;
53             returnValue = false;
54         }
55     }
56 
57     if (kernel != 0)
58     {
59         if (!checkSuccess(clReleaseKernel(kernel)))
60         {
61             cerr << "Releasing the OpenCL kernel failed. " << __FILE__ <<
62                  ":"<< __LINE__ << endl;
63             returnValue = false;
64         }
65     }
66 
67     if (program != 0)
68     {
69         if (!checkSuccess(clReleaseProgram(program)))
70         {
71             cerr << "Releasing the OpenCL program failed. " << __FILE__ <<
72                  ":"<< __LINE__ << endl;
73             returnValue = false;
74         }
75     }
76 
77     for (int index = 0; index < numberOfMemoryObjects; index++)
78     {
79         if (memoryObjects[index] != 0)
80         {
81             if (!checkSuccess(clReleaseMemObject(memoryObjects[index])))
82             {
83                 cerr << "Releasing the OpenCL memory object " << index <<
84                      " failed. " << __FILE__ <<
85                      ":"<< __LINE__ << endl;
86                 returnValue = false;
87             }
88         }
89     }
90 
91     return returnValue;
92 }
93 
94 
95 void
96 Cvt_2_Gray(Mat& src, Mat& dst)
97 {
98     int c_cannel = src.channels();
99     if ( c_cannel == 3 )
100         cvtColor(src, dst, CV_BGR2GRAY);
101     else if (c_cannel == 4)
102         cvtColor(src, dst, CV_BGRA2GRAY);
103     else if (c_cannel == 1)
104         src = dst;
105 }
106 
107 
108 cl_program
109 load_program(cl_context context,  string filename, cl_device_id device)
110 {
111     ifstream kernelFile(filename.c_str(), ios::in);
112 
113     if (!kernelFile.is_open())
114     {
115         cerr << "Unable to open " << filename << ". " << __FILE__ <<
116              ":"<< __LINE__ << endl;
117         return false;
118     }
119     ostringstream outputStringStream;
120     outputStringStream << kernelFile.rdbuf();
121     string srcStdStr = outputStringStream.str();
122     const char* source = srcStdStr.c_str();
123 
124 
125     cl_program program = clCreateProgramWithSource(context, 1, &source, 0, 0);
126     if (program == 0)
127     {
128         return 0;
129     }
130 
131     if (clBuildProgram(program, 0, 0, 0, 0, 0) != CL_SUCCESS)
132     {
133         char log[4096];
134         size_t log_size;
135         clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
136                               4096, log, &log_size);
137         cout << log << endl;
138         return 0;
139     }
140 
141     return program;
142 }
143 
144 
145 int main(void)
146 {
147     cl_context CPU_context = 0;
148     cl_context GPU_context = 0;
149     cl_command_queue commandQueue = 0;
150     cl_program program = 0;
151     cl_device_id device = 0;
152     cl_kernel kernel = 0;
153     int numberOfMemoryObjects = 2;
154     cl_mem memoryObjects[3] = {0, 0, 0};
155     cl_mem imgObject[3] = {0, 0, 0};
156     cl_int errorNumber;
157 
158     LARGE_INTEGER t1, t2, ts;
159     LARGE_INTEGER t1_tmp, t2_tmp;
160     double       CPU_elapse, GPU_elapse;
161     QueryPerformanceFrequency(&ts);
162 
163 
164     // OpenCV part.
165     Mat gray_Img0, gray_Img1;
166     Mat CPU_img;
167     Mat Img1 = cv::imread("0000.jpg");
168     if ( Img1.empty())
169         return 0;
170     Mat Img0;
171     cvtColor(Img1, Img0, CV_RGB2RGBA);
172     int ch = Img0.channels();
173 
174     imshow( "original image\n ", Img0 );
175     printf("img_w=%d, img_h=%d\n", Img0.cols, Img0.rows);
176 
177 
178     int kernel_size = 3;
179     Mat filter_kernel(kernel_size, kernel_size, CV_8S);
180     Point anchor = Point(-1,-1);
181     double delta;
182     int ddepth;
183     filter_kernel.data[0] = -1;
184     filter_kernel.data[1] = -1;
185     filter_kernel.data[2] = -1;
186     filter_kernel.data[3] = -1;
187     filter_kernel.data[4] = 9;
188     filter_kernel.data[5] = -1;
189     filter_kernel.data[6] = -1;
190     filter_kernel.data[7] = -1;
191     filter_kernel.data[8] = -1;
192     QueryPerformanceCounter(&t1);
193     filter2D( Img0, CPU_img, -1, filter_kernel, Point(-1,-1) );
194     QueryPerformanceCounter(&t2);
195     CPU_elapse = (double)(t2.QuadPart-t1.QuadPart)/(double)(ts.QuadPart);
196     printf("CPU Time consume: %lf\n", CPU_elapse );
197     imshow( "CPU image", CPU_img );
198 
199 
200 
201     cl_int err;
202     cl_uint num = 0;
203     cl_platform_id platforms = 0;
204 
205     if (!checkSuccess(clGetPlatformIDs(1, &platforms, &num)))
206     {
207         cerr << "Retrieving OpenCL platforms failed. " << __FILE__ <<
208              ":"<< __LINE__ << endl;
209         return false;
210     }
211 
212     cl_context_properties prop[]
213         = { CL_CONTEXT_PLATFORM,
214             reinterpret_cast<cl_context_properties>(platforms), 0
215           };
216     GPU_context = clCreateContextFromType(prop, CL_DEVICE_TYPE_GPU,
217                                           NULL, NULL, NULL);
218     if (GPU_context == 0)
219     {
220         cerr << "Can't create OpenCL context\n";
221         return 0;
222     }
223 
224     size_t deviceBUfSize;
225     clGetContextInfo(GPU_context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBUfSize);
226     vector<cl_device_id> devices(deviceBUfSize / sizeof(cl_device_id));
227     clGetContextInfo(GPU_context, CL_CONTEXT_DEVICES, deviceBUfSize,
228                      &devices[0], 0);
229 
230     // Get Device name
231     clGetDeviceInfo(devices[0], CL_DEVICE_NAME, 0, NULL, &deviceBUfSize);
232     string device_name;
233     device_name.resize(deviceBUfSize);
234     clGetDeviceInfo(devices[0], CL_DEVICE_NAME, deviceBUfSize, &device_name[0], 0);
235     std::cout << "Device: " << device_name.c_str() << "\n";
236 
237     // get device compute unit.
238     cl_uint compute_unit;
239     clGetDeviceInfo(devices[0], CL_DEVICE_MAX_COMPUTE_UNITS,
240                     sizeof(cl_uint), &compute_unit, 0);
241     cout << "Compute Unit: " << compute_unit << "\n";
242 
243 
244     // get max. work group size.
245     cl_uint max_work_group_size;
246     clGetDeviceInfo(devices[0], CL_DEVICE_MAX_WORK_GROUP_SIZE,
247                     sizeof(cl_uint), &max_work_group_size, 0);
248 
249     // get some infomation:
250     cl_uint tmp;
251     clGetDeviceInfo(devices[0], CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR,
252                     sizeof(cl_uint), &tmp, 0);
253     clGetDeviceInfo(devices[0], CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT,
254                     sizeof(cl_uint), &tmp, 0);
255     clGetDeviceInfo(devices[0], CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT,
256                     sizeof(cl_uint), &tmp, 0);
257     clGetDeviceInfo(devices[0], CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE,
258                     sizeof(cl_uint), &tmp, 0);
259 
260     clGetDeviceInfo(devices[0], CL_DEVICE_MAX_CLOCK_FREQUENCY,
261                     sizeof(cl_uint), &tmp, 0);
262 
263     // check if device support image
264     cl_bool img_support;
265     clGetDeviceInfo(devices[0], CL_DEVICE_IMAGE_SUPPORT,
266                     sizeof(cl_bool), &img_support, 0);
267     if ( img_support == CL_TRUE )
268         printf("device support image\n");
269 
270     // create command queue:
271     commandQueue = clCreateCommandQueue(GPU_context, devices[0], 0, 0);
272     if (commandQueue == 0)
273     {
274         std::cerr << "Can't create command queue\n";
275         clReleaseContext(GPU_context);
276         return 0;
277     }
278     //int w = Img0.Width;
279     // generate testing data.
280     const int DATA_SIZE = Img0.rows * Img0.cols * Img0.channels();
281     std::vector<cl_uchar> res(DATA_SIZE);
282     //char *buffer = new char [Img0.rows * Img0.cols * Img0.channels()];
283 
284     Mat output_img(Img0.rows, Img0.cols,CV_8UC4, Scalar::all(0) );
285 
286     // create image object
287     cl_image_format input_img_fmt;
288     input_img_fmt.image_channel_order = CL_RGBA;
289     input_img_fmt.image_channel_data_type = CL_UNSIGNED_INT8;
290 
291     cl_image_format output_img_fmt;
292     output_img_fmt.image_channel_order = CL_RGBA;
293     output_img_fmt.image_channel_data_type = CL_UNSIGNED_INT8;
294 
295     // create image object
296     imgObject[0] =
297         clCreateImage2D(GPU_context,
298                         CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
299                         &input_img_fmt, Img0.cols, Img0.rows,
300                         0, Img0.data, NULL);
301     imgObject[1] =
302         clCreateImage2D(GPU_context, CL_MEM_WRITE_ONLY,
303                         &output_img_fmt, Img0.cols, Img0.rows,
304                         0, NULL, NULL);
305 
306 
307     // load program and build.
308     char shader_name[] = "filter_shader_img.cl";
309     printf("%s\n", shader_name);
310     program = load_program( GPU_context, shader_name, devices[0] );
311     if ( program == 0 )
312     {
313         cleanUpOpenCL(GPU_context, commandQueue, program, kernel,
314                       memoryObjects, numberOfMemoryObjects);
315         return 0;
316     }
317 
318     kernel = clCreateKernel( program, "laplacian", NULL );
319     if ( kernel == 0 )
320     {
321         cleanUpOpenCL(GPU_context, commandQueue, program, kernel,
322                       memoryObjects, numberOfMemoryObjects);
323         return 0;
324     }
325 
326     // config kernel argument
327     // source image
328     clSetKernelArg(kernel, 0, sizeof(cl_mem), &imgObject[0]);
329     // output image
330     clSetKernelArg(kernel, 1, sizeof(cl_mem), &imgObject[1]);
331 
332     const int workDimensions = 2;
333     size_t globalWorkSize[workDimensions] = {Img0.cols, Img0.rows};
334 
335     cl_event evt;
336     cl_event r_evt;
337     err = clEnqueueNDRangeKernel(commandQueue, kernel, workDimensions,
338                                  NULL,globalWorkSize, NULL,
339                                  0, NULL, &evt);
340     //clFinish( commandQueue );
341     if (err == CL_SUCCESS)
342     {
343         size_t origin[3] = {0,0,0};
344         size_t region[3] = {Img0.cols, Img0.rows, 1};
345         err = clEnqueueReadImage(commandQueue, imgObject[1], CL_FALSE,
346                                  origin, region, 0, 0, &res[0],
347                                  1, &evt, &r_evt);
348     }
349 
350     clWaitForEvents (1, &r_evt);
351     Mat ocl_img(Size(Img0.cols, Img0.rows), CV_8UC4, &res[0]);
352 
353     imshow("GPU image", ocl_img );
354 
355 
356 
357     // Release OpenCL objects.
358     res.clear();
359     cleanUpOpenCL(GPU_context, commandQueue, program, kernel,
360                   imgObjects, numberOfMemoryObjects);
361     cvWaitKey(0);
362 }
363