現在，我們開始寫一個簡單的OpenCL程式，計算兩個數組相加的和，放到另一個數組中去。程式用cpu和gpu分别計算，最後驗證它們是否相等。OpenCL程式的流程大緻如下：
下面是source code中的主要代碼：
int main(int argc, char* argv[])
{
//在host記憶體中建立三個緩沖區
float *buf1 = 0;
float *buf2 = 0;
float *buf = 0;
buf1 =(float *)malloc(BUFSIZE * sizeof(float));
buf2 =(float *)malloc(BUFSIZE * sizeof(float));
buf =(float *)malloc(BUFSIZE * sizeof(float));
//用一些随機值初始化buf1和buf2的内容
int i;
srand( (unsigned)time( NULL ) );
for(i = 0; i < BUFSIZE; i++)
buf1[i] = rand()%65535;
srand( (unsigned)time( NULL ) +1000);
buf2[i] = rand()%65535;
//cpu計算buf1,buf2的和
buf[i] = buf1[i] + buf2[i];
cl_uint status;
cl_platform_id platform;
//建立平台對象
status = clGetPlatformIDs( 1, &platform, NULL );
注意：如果我們系統中安裝不止一個opencl平台，比如我的os中，有intel和amd兩家opencl平台，用上面這行代碼，有可能會出錯，因為它得到了intel的opencl平台，而intel的平台隻支援cpu，而我們後面的操作都是基于gpu，這時我們可以用下面的代碼，得到AMD的opencl平台。
cl_device_id device;
//建立GPU裝置
clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
1,
&device,
NULL);
//建立context
cl_context context = clCreateContext( NULL,
NULL, NULL, NULL);
//建立指令隊列
cl_command_queue queue = clCreateCommandQueue( context,
device,
CL_QUEUE_PROFILING_ENABLE, NULL );
//建立三個OpenCL記憶體對象，并把buf1的内容通過隐式拷貝的方式
//buf1内容拷貝到clbuf1,buf2的内容通過顯示拷貝的方式拷貝到clbuf2
cl_mem clbuf1 = clCreateBuffer(context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
BUFSIZE*sizeof(cl_float),buf1,
NULL );
cl_mem clbuf2 = clCreateBuffer(context,
CL_MEM_READ_ONLY ,
BUFSIZE*sizeof(cl_float),NULL,
cl_event writeEvt;
status = clEnqueueWriteBuffer(queue, clbuf2, 1,
0, BUFSIZE*sizeof(cl_float), buf2, 0, 0, 0);
上面這行代碼把buf2中的内容拷貝到clbuf2,因為buf2位于host端，clbuf2位于device端，是以這個函數會執行一次host到device的傳輸操作，或者說一次system memory到video memory的拷貝操作，是以我在該函數的後面放置了clFush函數，表示把command queue中的所有指令送出到device(注意：該指令并不保證指令執行完成),是以我們調用函數waitForEventAndRelease來等待write緩沖的完成，waitForEventAndReleae 是一個使用者定義的函數，它的内容如下，主要代碼就是通過event來查詢我們的操作是否完成，沒完成的話，程式就一直block在這行代碼處，另外我們也可以用opencl中内置的函數clWaitForEvents來代替clFlush和waitForEventAndReleae。
status = clFlush(queue);
//等待資料傳輸完成再繼續往下執行
waitForEventAndRelease(&writeEvt);
cl_mem buffer = clCreateBuffer( context,
CL_MEM_WRITE_ONLY,
BUFSIZE * sizeof(cl_float),
NULL, NULL );
kernel檔案中放的是gpu中執行的代碼，它被放在一個單獨的檔案add.cl中，本程式中kernel代碼非常簡單，隻是執行兩個數組相加。kernel的代碼為：
//kernel檔案為add.cl
const char * filename = "add.cl";
std::string sourceStr;
status = convertToString(filename, sourceStr);
convertToString也是使用者定義的函數，該函數把kernel源檔案讀入到一個string中，它的代碼如下：
const char * source = sourceStr.c_str();
size_t sourceSize[] = { strlen(source) };
//建立程式對象
cl_program program = clCreateProgramWithSource(
context,
&source,
sourceSize,
//編譯程式對象
status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );
if(status != 0)
{
printf("clBuild failed:%d\n", status);
char tbuf[0x10000];
clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);
printf("\n%s\n", tbuf);
return -1;
}
//建立Kernel對象
cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );
//設定Kernel參數
cl_int clnum = BUFSIZE;
clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);
clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);
clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);
注意：在執行kernel時候，我們隻設定了global work items數量，沒有設定group size，這時候，系統會使用預設的work group size，通常可能是256之類的。
//執行kernel,Range用1維，work itmes size為BUFSIZE
cl_event ev;
size_t global_work_size = BUFSIZE;
clEnqueueNDRangeKernel( queue,
kernel,
NULL,
&global_work_size,
NULL, 0, NULL, &ev);
status = clFlush( queue );
waitForEventAndRelease(&ev);
//資料拷回host記憶體
cl_float *ptr;
cl_event mapevt;
ptr = (cl_float *) clEnqueueMapBuffer( queue,
buffer,
CL_TRUE,
CL_MAP_READ,
0,
0, NULL, NULL, NULL );
waitForEventAndRelease(&mapevt);
//結果驗證，和cpu計算的結果比較
if(!memcmp(buf, ptr, BUFSIZE))
printf("Verify passed\n");
else printf("verify failed");
if(buf)
free(buf);
if(buf1)
free(buf1);
if(buf2)
free(buf2);
程式結束後，這些opencl對象一般會自動釋放，但是為了程式完整，養成一個好習慣，這兒我加上了手動釋放opencl對象的代碼。
//删除OpenCL資源對象
clReleaseMemObject(clbuf1);
clReleaseMemObject(clbuf2);
clReleaseMemObject(buffer);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}
程式執行後的界面如下：
完整的代碼請參考：
工程檔案gclTutorial1
代碼下載下傳：
<a href="http://files.cnblogs.com/mikewolf2002/gclTutorial.zip">http://files.cnblogs.com/mikewolf2002/gclTutorial.zip</a>