原文http://www.olcf.ornl.gov/training_articles/opencl-vector-addition/

本文僅僅是為了學(xué)習(xí)OpenCL而做的的相關(guān)翻譯。

由于原文中的例子不能在我的環(huán)境中運(yùn)行，因此做了一些改動(dòng)。

通過這個(gè)例子能很好地了解OpenCL的編程模型。

1. 簡介

這個(gè)例子是表示了兩個(gè)向量相加，可以認(rèn)為是OpenCL中的"hello world"。為了使程序更容易理解，沒有加入錯(cuò)誤處理機(jī)制。

//vecAdd.c#include <stdio.h> #include <stdlib.h> #include <math.h> #include <CL/opencl.h>// OpenCL kernel. Each work item takes care of one element of c const char *kernelSource = "\n" \ "__kernel void vecAdd( __global float *a, \n" \ " __global float *b, \n" \ " __global float *c, \n" \ " const unsigned int n) \n" \ "{ \n" \ " //Get our global thread ID \n" \ " int id = get_global_id(0); \n" \ " \n" \ " //Make sure we do not go out of bounds \n" \ " if (id < n) \n" \ " c[id] = a[id] + b[id]; \n" \ "} \n" \"\n" ;int main( int argc, char* argv[] ) {// 向量長度int n = 8;// 輸入向量int *h_a;int *h_b;// 輸出向量int *h_c;// 設(shè)備輸入緩沖區(qū)cl_mem d_a;cl_mem d_b;// 設(shè)備輸出緩沖區(qū)cl_mem d_c;cl_platform_id cpPlatform; // OpenCL 平臺(tái)cl_device_id device_id; // device IDcl_context context; // contextcl_command_queue queue; // command queuecl_program program; // programcl_kernel kernel; // kernel//（每個(gè)向量的字節(jié)數(shù)）size_t bytes = n*sizeof(int);//（為每個(gè)向量分配內(nèi)存）h_a = (int*)malloc(bytes);h_b = (int*)malloc(bytes);h_c = (int*)malloc(bytes);//（初始化向量）int i;for( i = 0; i < n; i++ ){h_a[i] = i;h_b[i] = i;}size_t globalSize, localSize;cl_int err;//（每個(gè)工作組的工作節(jié)點(diǎn)數(shù)目）localSize = 2;//（所有的工作節(jié)點(diǎn)）globalSize = (size_t)ceil(n/(float)localSize)*localSize;printf("%d\n",globalSize);//（獲得平臺(tái)ID）err = clGetPlatformIDs(1, &cpPlatform, NULL);//（獲得設(shè)備ID，與平臺(tái)有關(guān)）err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);//（根據(jù)設(shè)備ID，得到上下文）context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);//（根據(jù)上下文，在設(shè)備上創(chuàng)建命令隊(duì)列）queue = clCreateCommandQueue(context, device_id, 0, &err);//（根據(jù)OpenCL源程序創(chuàng)建計(jì)算程序）program = clCreateProgramWithSource(context, 1,(const char **) & kernelSource, NULL, &err);//（創(chuàng)建可執(zhí)行程序）clBuildProgram(program, 0, NULL, NULL, NULL, NULL);//（在上面創(chuàng)建的程序中創(chuàng)建內(nèi)核程序）kernel = clCreateKernel(program, "vecAdd", &err);//（分配設(shè)備緩沖）d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);// （將向量信息寫入設(shè)備緩沖）err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,bytes, h_a, 0, NULL, NULL);err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,bytes, h_b, 0, NULL, NULL);// （設(shè)置計(jì)算內(nèi)核的參數(shù)）err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);err |= clSetKernelArg(kernel, 3, sizeof(int), &n);// （在數(shù)據(jù)集的范圍內(nèi)執(zhí)行內(nèi)核）Execute the kernel over the entire range of the data seterr = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,0, NULL, NULL);// （在讀出結(jié)果之前，等待命令隊(duì)列執(zhí)行完畢）Wait for the command queue to get serviced before reading back resultsclFinish(queue);// （從設(shè)備緩沖區(qū)讀出結(jié)果）Read the results from the deviceclEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,bytes, h_c, 0, NULL, NULL );//（輸出讀出的結(jié)果）float sum = 0;for(i=0; i<n; i++)printf("%d ",h_c[i]);// （釋放資源）clReleaseMemObject(d_a);clReleaseMemObject(d_b);clReleaseMemObject(d_c);clReleaseProgram(program);clReleaseKernel(kernel);clReleaseCommandQueue(queue);clReleaseContext(context);//（釋放內(nèi)存）free(h_a);free(h_b);free(h_c);system("pause");return 0; }

2. 基本解釋

2.1 內(nèi)核：

Kernel是OpenCL代碼的核心。全部kernel必須要作為一個(gè)C字符串讀入，最容易的方式就是將整個(gè)kernel用引號(hào)包起來，行尾回車。在真正的程序中，應(yīng)該將kernel放在一個(gè)獨(dú)立的文件中。

// OpenCL kernel. Each work item takes care of one element of c const char *kernelSource = "\n" \ "__kernel void vecAdd( __global float *a, \n" \ " __global float *b, \n" \ " __global float *c, \n" \ " const unsigned int n) \n" \ "{ \n" \ " //Get our global thread ID \n" \ " int id = get_global_id(0); \n" \ " \n" \ " //Make sure we do not go out of bounds \n" \ " if (id < n) \n" \ " c[id] = a[id] + b[id]; \n" \ "} \n" \"\n" ;

　　

查看一下這個(gè)簡單的內(nèi)核由什么內(nèi)容組成：

__kernel void vecAdd( __global float *a, __global float *b,
__global float *c, const unsigned int n)

__kernel 指明這是一個(gè)OpenCL內(nèi)核，__global 說明指針指向的是全局的設(shè)備內(nèi)存空間，其它的就是C語言的函數(shù)的語法。kernel必須返回空類型。

int id = get_global_id(0);

得到第0維全局工作節(jié)點(diǎn)的ID。

if (id < n)
c[id] = a[id] + b[id];

工作組的數(shù)目必須是一個(gè)整數(shù)，或者每個(gè)工作組的工作節(jié)點(diǎn)數(shù)目必須能被全部工作節(jié)點(diǎn)數(shù)目整除。由于共組的的大小被用來協(xié)調(diào)性能，沒有必要一定能被所有線程數(shù)目整除，所以通常啟用的線程比所需要的線程多一些，并忽略掉多余的。在考察了問題域之后，就能訪問、操作設(shè)備內(nèi)存了。
2.2?內(nèi)存：

// 輸入向量
int *h_a;
int *h_b;
// 輸出向量
int *h_c;

// 設(shè)備輸入緩沖區(qū)
cl_mem d_a;
cl_mem d_b;
// 設(shè)備輸出緩沖區(qū)
cl_mem d_c;

CPU和GPU有不同的內(nèi)存空間，所以必須支持對(duì)內(nèi)存分別引用，一個(gè)集市主機(jī)數(shù)組指針，另外一個(gè)集是設(shè)備內(nèi)存的操作句柄。這兒我們用 h_和d_前綴來區(qū)分。

2.3 線程映射：

//（每個(gè)工作組的工作節(jié)點(diǎn)數(shù)目）
localSize = 2;

//（所有的工作節(jié)點(diǎn)）
globalSize = (size_t)ceil(n/(float)localSize)*localSize;

為了將問題映射到底層硬件，必須指明局部的大小，和全局的大小。局部大小定義了工作組中節(jié)點(diǎn)的數(shù)目，子NVIDIA GPU上這相當(dāng)于線程塊內(nèi)線程的數(shù)目。全局大小定義了所有啟動(dòng)的工作節(jié)點(diǎn)數(shù)目。localSize大小必須能被globalSize整除，所以我們計(jì)算了一個(gè)最小的整數(shù)能覆蓋問題域，并且能被localSize整除。

2.4 環(huán)境配置：

?

//（綁定平臺(tái)）
err = clGetPlatformIDs(1, &cpPlatform, NULL);

每個(gè)硬件提供商都有不同的平臺(tái)，在用之前就應(yīng)該給定，這兒clGetPlatformIDs()會(huì)將cpPlatform賦予系統(tǒng)可用的平臺(tái)。例如，如果系統(tǒng)包含了AMD CPU和NVIDIA GPU，且這兩個(gè)平臺(tái)都安裝了合適的OpenCL驅(qū)動(dòng)，那這里平臺(tái)都是可用的。（注：要使用不同的平臺(tái)驅(qū)動(dòng)，必須安裝相關(guān)的驅(qū)動(dòng)，在本例中我安裝了AMD(ATI)的app SDK v2.5和Intel的intel_ocl_sdk_1.5_runtime_setup，所以會(huì)有兩個(gè)平臺(tái)，但是由于我的ATI的顯卡GPU不能被app SDK v2.5支持，所以的獲得設(shè)備ID時(shí)沒有用GPU設(shè)備，而是用了CPU設(shè)備。如果這里配置不正確，下面的可能就無法進(jìn)行）

?

//（獲得設(shè)備ID，與平臺(tái)有關(guān)）
err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);

可以查詢平臺(tái)來得到它包含什么樣的設(shè)備。在這個(gè)例子中，用枚舉值CL_DEVICE_TYPE_CPU來查詢平臺(tái)上的CPU設(shè)備。

?

//（根據(jù)上下文，在設(shè)備上創(chuàng)建命令隊(duì)列）
queue = clCreateCommandQueue(context, device_id, 0, &err);

在使用OpenCL設(shè)備之前，必須要配置context，context被用來管理命令隊(duì)列，內(nèi)存和內(nèi)核的活動(dòng)。一個(gè)context可以包含不止一個(gè)設(shè)備。

//（根據(jù)上下文，在設(shè)備上創(chuàng)建命令隊(duì)列）
queue = clCreateCommandQueue(context, device_id, 0, &err);

命令隊(duì)列被用來將命令從主機(jī)放入指定的設(shè)備。內(nèi)存的轉(zhuǎn)移和內(nèi)核的活動(dòng)都能被放入命令隊(duì)列在合適的時(shí)候在指定的設(shè)備上執(zhí)行。

2.5 編譯內(nèi)核：

//（根據(jù)OpenCL源程序創(chuàng)建計(jì)算程序）
program = clCreateProgramWithSource(context, 1,
(const char **) & kernelSource, NULL, &err);

//（創(chuàng)建可執(zhí)行程序）
clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

//（在上面創(chuàng)建的程序中創(chuàng)建內(nèi)核程序）
kernel = clCreateKernel(program, "vecAdd", &err);

為了保證對(duì)于大多數(shù)設(shè)備的可移植性，默認(rèn)運(yùn)行內(nèi)核的方式就是用即時(shí)（Just-in-time）編譯我們必須為給定上下文的設(shè)備準(zhǔn)備源碼。首先，創(chuàng)建程序，這是一個(gè)內(nèi)核程序的的集合，然后根據(jù)程序來創(chuàng)建各自的內(nèi)核程序。

2.6 準(zhǔn)備數(shù)據(jù)：

?

//（分配設(shè)備緩沖）
d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);
d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);

// （將向量信息寫入設(shè)備緩沖）
err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,
bytes, h_a, 0, NULL, NULL);
err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,
bytes, h_b, 0, NULL, NULL);

// （設(shè)置計(jì)算內(nèi)核的參數(shù)）
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);
err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);
err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);
err |= clSetKernelArg(kernel, 3, sizeof(int), &n);

在啟動(dòng)內(nèi)核之前，必須在設(shè)備和主機(jī)之間創(chuàng)建緩沖區(qū)，將主機(jī)數(shù)據(jù)綁定到新創(chuàng)建的設(shè)備緩沖區(qū)上，最后，設(shè)置內(nèi)核參數(shù)。

2.7 啟動(dòng)內(nèi)核：

?

// （在數(shù)據(jù)集的范圍內(nèi)執(zhí)行內(nèi)核）
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,
0, NULL, NULL);

一旦內(nèi)存駐留在設(shè)備上之后，內(nèi)核就能排隊(duì)啟動(dòng)了。

2.8 取回結(jié)果：

?

// （在讀出結(jié)果之前，等待命令隊(duì)列執(zhí)行完畢）
clFinish(queue);

// （從設(shè)備緩沖區(qū)讀出結(jié)果）
clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,
bytes, h_c, 0, NULL, NULL );

可以進(jìn)行阻斷，直到所有的命令隊(duì)列執(zhí)行完畢，然后將設(shè)備上的結(jié)果取回到主機(jī)。

3. 運(yùn)行環(huán)境

3.1 ?OpenCL:

　　AMD app sdk v2.5

　　intel_ocl_sdk_1.5_runtime

3.2 Visual Studio 2010 express

轉(zhuǎn)載于:https://www.cnblogs.com/wangshide/archive/2011/11/04/2235204.html

總結(jié)

以上是生活随笔為你收集整理的OpenCL向量相加的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯(cuò)，歡迎將生活随笔推薦給好友。