AMD OpenCL大学教程(8)

在本节，我们主要介绍OpenCL中buffer的使用，同时提供了两个个完整的例子，一个是图像的旋转，一个是矩阵乘法（非常简单，没有分块优化）。

1、创建OpenCL设备缓冲(buffer)

   OpenCL设备使用的数据都存放在设备的buffer中[其实就是device memory中]。我们用下面的代码创建buffer对象：

cl_mem bufferobj = clCreateBuffer ( 
cl_context context, //Context name
cl_mem_flags flags, //Memory flags
size_t size, //Memory size allocated in buffer
 void *host_ptr, //Host data 
cl_int *errcode) //Returned error code

    如果host_ptr指向一个有效的host指针，则创建一个buffer对象的同时会实现隐式的数据拷贝（会在kernel函数进入队列时候，把host_prt中的数据从host memory拷贝到设备内存对象bufferobj中)。

    我们可以通过flags参数指定buffer对象的属性。

   函数clEnqueueWriteBuffer()用来实现显示的数据拷贝，即把host memory中的数据拷贝到device meomory中。

cl_int clEnqueueWriteBuffer ( 
cl_command_queue queue, //Command queue to device
cl_mem buffer, //OpenCL Buffer Object
cl_bool blocking_read, //Blocking/Non-Blocking Flag
size_t offset, //Offset into buffer to write to
size_t cb, //Size of data
 void *ptr, //Host pointer
cl_uint num_in_wait_list, //Number of events in wait list
 const cl_event * event_wait_list, //Array of events to wait for
cl_event *event) //Event handler for this function

2、图像旋转的例子

   下面是一个完整的OpenCL例子，实现图像的旋转。在这个例子中，我把美丽的Lenna旋转了90度。

下面是原始图像和旋转后的图像（黑白）

在这个例子中，我使用FreeImage库，可以从FreeImage网站或者我的code工程中下载。

http://code.google.com/p/imagefilter-opencl/downloads/detail?name=Dist.rar&can=2&q=#makechanges

   图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度，通常是指绕图像的中心以逆时针方向旋转。

假设图像的左上角为（left, top),右下角为（right, bottom)，则图像上任意点（x0, y0）绕其中心（xcenter, ycenter)逆时针旋转angle角度后，新的坐标位置（x′, y′）的计算公式为：

xcenter = (right － left ＋ 1) ／ 2 ＋ left;
ycenter = (bottom － top ＋ 1) ／ 2 ＋ top;
x′ = (x0 － xcenter) cosθ － (y0 － ycenter) sinθ ＋ xcenter;
y′ = (x0 － xcenter) sinθ ＋ (y0 － ycenter) cosθ ＋ ycenter;

下面给出kernel的代码：

1: __kernel void image_rotate( __global uchar * src_data, __global uchar * dest_data, //Data in global memory
2: int W, int H, //Image Dimensions
3: float sinTheta, float cosTheta ) //Rotation Parameters
4: {
5: //Thread gets its index within index space
6: const int ix = get_global_id(0);
7: const int iy = get_global_id(1);
8:
9: int xc = W/2;
10: int yc = H/2;
11:
12: int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;
13: int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;
14:
15: if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H)) //Bound Checking
16: {
17: dest_data[ypos*W+xpos]= src_data[iy*W+ix];
18: }
19: }
20:

src_data为原始图像（灰度图）数据，dest_data为旋转后的图像数据。W、H分别为图像的高度和宽度。sinTheta和cosTheta是旋转参数。我在代码中实现了旋转90度，所以sinTheta为1，cosTheta为0，大家可以尝试其它的值。

下面是程序的流程图：

在前面向量加法的例子中，我已经介绍了OpenCL一些基本的步骤。

1. 创建platform对象
2. 创建GPU设备
3. 创建contex
4. 创建命令队列
5. 创建缓冲对象，代码如下：

   1: cl_mem d_ip = clCreateBuffer(
   2: context, CL_MEM_READ_ONLY,
   3: mem_size,
   4: NULL, NULL);
   5: l_mem d_op = clCreateBuffer(
   6: context, CL_MEM_WRITE_ONLY,
   7: mem_size,
   8: NULL, NULL);
   9: status = clEnqueueWriteBuffer (
   10: queue , d_ip, CL_TRUE,
   11: 0, mem_size, (void *)src_image,
   12: 0, NULL, NULL);

6. 创建程序对象
7. 编译程序对象
8. 创建Kernel对象
9. 设置kernel参数
10. 执行kernel
11. 数据拷贝回host memory，我采用映射memory的方式

1: unsigned char *op_data=0;
2: //op_data =(unsigned char *)malloc(mem_size);
3: status = clEnqueueReadBuffer(
4: //queue, d_op,
5: //CL_TRUE, //Blocking Read Back
6: //0, mem_size,(void*)op_data, NULL, NULL, NULL);
7: op_data = (cl_uchar *) clEnqueueMapBuffer( queue,
8: d_op,
9: CL_TRUE,
10: CL_MAP_READ,
11: 0,
12: mem_size,
13: 0, NULL, NULL, NULL );

   kernel执行时间的计算后面教程会有详细介绍，但在本节中，我们会给出通过事件机制来得到kernel执行时间，首先要在创建队列时候，使用CL_QUEUE_PROFILING_ENABLE参数，否则计算的kernel运行时间是0。

   下面是代码：

 1: //计算kerenl执行时间
 2: cl_ulong startTime, endTime;
clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START,
 4: sizeof(cl_ulong), &startTime, NULL);
 5: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END,
 6: sizeof(cl_ulong), &endTime, NULL);
 7: cl_ulong kernelExecTimeNs = endTime-startTime;
 8: printf("kernal exec time :%8.6f ms\n ", kernelExecTimeNs*1e-6 );

完整的程序代码：

1: #include "stdafx.h"
2: #include <CL/cl.h>
3: #include <stdio.h>
4: #include <stdlib.h>
5: #include <time.h>
6: #include <iostream>
7: #include <fstream>
8:
9: #include "gFreeImage.h"
10:
11: using namespace std;
12: #define NWITEMS 4
13: #pragma comment (lib,"OpenCL.lib")
14: #pragma comment(lib,"FreeImage.lib")
15:
16: //把文本文件读入一个string中
17: int convertToString(const char *filename, std::string& s)
18: {
19: size_t size;
20: char* str;
21:
22: std::fstream f(filename, (std::fstream::in | std::fstream::binary));
23:
24: if(f.is_open())
25: {
26: size_t fileSize;
27: f.seekg(0, std::fstream::end);
28: size = fileSize = (size_t)f.tellg();
29: f.seekg(0, std::fstream::beg);
30:
31: str = new char[size+1];
32: if(!str)
33: {
34: f.close();
35: return NULL;
36: }
37:
38: f.read(str, fileSize);
39: f.close();
40: str[size] = '\0';
41:
42: s = str;
43: delete[] str;
44: return 0;
45: }
46: printf("Error: Failed to open file %s\n", filename);
47: return 1;
48: }
49:
50: //CPU旋转图像
51: void cpu_rotate(unsigned char* inbuf, unsigned char* outbuf, int w, int h,float sinTheta, float cosTheta)
52: {
53: int i, j;
54: int xc = w/2;
55: int yc = h/2;
56:
57: for(i = 0; i < h; i++)
58: {
59: for(j=0; j< w; j++)
60: {
61: int xpos = ( j-xc)*cosTheta - (i-yc)*sinTheta+xc;
62: int ypos = (j-xc)*sinTheta + ( i-yc)*cosTheta+yc;
63:
64: if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
65: outbuf[ypos*w + xpos] = inbuf[i*w+j];
66: }
67: }
68: }
69:
70: int main(int argc, char* argv[])
71: {
72: //装入图像
73: unsigned char *src_image=0;
74: unsigned char *cpu_image=0;
75: int W, H;
76: gFreeImage img;
77: if(!img.LoadImageGrey("lenna.jpg"))
78: {
79: printf("装入lenna.jpg失败\n");
80: exit(0);
81: }
82: else
83: src_image = img.getImageDataGrey(W, H);
84:
85: size_t mem_size = W*H;
86: cpu_image = (unsigned char*)malloc(mem_size);
87:
88: cl_uint status;
89: cl_platform_id platform;
90:
91: //创建平台对象
92: status = clGetPlatformIDs( 1, &platform, NULL );
93:
94: cl_device_id device;
95:
96: //创建GPU设备
97: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
98: 1,
99: &device,
100: NULL);
101: //创建context
102: cl_context context = clCreateContext( NULL,
103: 1,
104: &device,
105: NULL, NULL, NULL);
106: //创建命令队列
107: cl_command_queue queue = clCreateCommandQueue( context,
108: device,
109: CL_QUEUE_PROFILING_ENABLE, NULL );
110:
111: //创建三个OpenCL内存对象，并把buf1的内容通过隐式拷贝的方式
112: //拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2
113: cl_mem d_ip = clCreateBuffer(
114: context, CL_MEM_READ_ONLY,
115: mem_size,
116: NULL, NULL);
117: cl_mem d_op = clCreateBuffer(
118: context, CL_MEM_WRITE_ONLY,
119: mem_size,
120: NULL, NULL);
121: status = clEnqueueWriteBuffer (
122: queue , d_ip, CL_TRUE,
123: 0, mem_size, (void *)src_image,
124: 0, NULL, NULL);
125:
126: const char * filename = "rotate.cl";
127: std::string sourceStr;
128: status = convertToString(filename, sourceStr);
129: const char * source = sourceStr.c_str();
130: size_t sourceSize[] = { strlen(source) };
131:
132: //创建程序对象
133: cl_program program = clCreateProgramWithSource(
134: context,
135: 1,
136: &source,
137: sourceSize,
138: NULL);
139: //编译程序对象
140: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );
141: if(status != 0)
142: {
143: printf("clBuild failed:%d\n", status);
144: char tbuf[0x10000];
145: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);
146: printf("\n%s\n", tbuf);
147: return -1;
148: }
149:
150:
151: //创建Kernel对象
152: //Use the “image_rotate” function as the kernel
153:
154: //创建Kernel对象
155: cl_kernel kernel = clCreateKernel( program, "image_rotate", NULL );
156:
157: //设置Kernel参数
158: float sintheta = 1, costheta = 0;
159: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&d_ip);
160: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&d_op);
161: clSetKernelArg(kernel, 2, sizeof(cl_int), (void *)&W);
162: clSetKernelArg(kernel, 3, sizeof(cl_int), (void *)&H);
163: clSetKernelArg(kernel, 4, sizeof(cl_float), (void *)&sintheta);
164: clSetKernelArg(kernel, 5, sizeof(cl_float), (void *)&costheta);
165:
166: //Set local and global workgroup sizes
167: size_t localws[2] = {16,16} ;
168: size_t globalws[2] = {W, H};//Assume divisible by 16
169:
170: cl_event ev;
171: //执行kernel
172: clEnqueueNDRangeKernel(
173: queue ,kernel,
174: 2, 0, globalws, localws,
175: 0, NULL, &ev);
176:
177: clFinish( queue );
178:
179: //计算kerenl执行时间
180: cl_ulong startTime, endTime;
181: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_START,
182: sizeof(cl_ulong), &startTime, NULL);
183: clGetEventProfilingInfo(ev, CL_PROFILING_COMMAND_END,
184: sizeof(cl_ulong), &endTime, NULL);
185: cl_ulong kernelExecTimeNs = endTime-startTime;
186: printf("kernal exec time :%8.6f ms\n ", kernelExecTimeNs*1e-6 );
187:
188: //数据拷回host内存
189: // copy results from device back to host
190: unsigned char *op_data=0;
191: //op_data =(unsigned char *)malloc(mem_size);
192: // status = clEnqueueReadBuffer(
193: //queue, d_op,
194: //CL_TRUE, //Blocking Read Back
195: //0, mem_size,(void*)op_data, NULL, NULL, NULL);
196: op_data = (cl_uchar *) clEnqueueMapBuffer( queue,
197: d_op,
198: CL_TRUE,
199: CL_MAP_READ,
200: 0,
201: mem_size,
202: 0, NULL, NULL, NULL );
203:
204: int i;
205: cpu_rotate(src_image,cpu_image, W, H, 1, 0);
206: for(i = 0; i < mem_size; i++)
207: {
208: src_image[i] =cpu_image[i];
209: }
210: img.SaveImage("cpu_lenna_rotate.jpg");
211: for(i = 0; i < mem_size; i++)
212: {
213: src_image[i] =op_data[i];
214: }
215: img.SaveImage("lenna_rotate.jpg");
216:
217: if(cpu_image)
218: free(cpu_image);
219:
220: //删除OpenCL资源对象
221: clReleaseMemObject(d_ip);
222: clReleaseMemObject(d_op);
223: clReleaseProgram(program);
224: clReleaseCommandQueue(queue);
225: clReleaseContext(context);
226: return 0;
227: }
228:

感兴趣的朋友可以从http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode2.zip&can=2&q=#makechanges下载完整代码。

注意代码运行后，会在程序目录生成lenna_rotate.jpg,这时gpu执行的结果，另外还有一个cpu_lenna_rotate.jpg这是CPU执行的结果。

3、一个矩阵乘法的例子

    在amd的slides中，本节还讲了一个简单的，没有优化的矩阵乘法，一共才两页ppt，所以我也不在这儿详细讲述了，…但简单介绍还是需要的。

1: for(int i = 0; i < Ha; i++)
2: for(int j = 0; j < Wb; j++){
3: c[i][j] = 0;
4: for(int k = 0; k < Wa; k++)
5: c[i][j] += a[i][k] + b[k][j]
6: }

上面的代码是矩阵乘法的例子，有三重循环，下面我们只给出kernel代码，完整程序请从：http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicodeCode3.zip&can=2&q=#makechanges下载。

1: __kernel void simpleMultiply(
2: __global float* c, int Wa, int Wb,
3: __global float* a, __global float* b)
4: {
5:
6: //Get global position in Y direction
7: int row = get_global_id(1);
8: //Get global position in X direction
9: int col = get_global_id(0);
10: float sum = 0.0f;
11: //Calculate result of one element
12: for (int i = 0; i < Wa; i++)
13: {
14: sum +=
15: a[row*Wa+i] * b[i*Wb+col];
16: }
17: c[row*Wb+col] = sum;
18: }
19:
20: