AMD OpenCL大学课程(5)

OpenCL内存模型

    OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的，比如从host memory到device memory，从global memory到local memory等等。

    WorkGroup被映射到硬件的CU上执行（在AMD 5xxx系列显卡上，CU就是simd，一个simd中有16个pe），OpenCL并不提供各个workgroup之间的一致性，如果我们需要在各个workgroup之间共享数据或者通信之类的，要自己通过软件实现。

Kernel函数的写法

每个线程（workitem）都有一个kenerl函数的实例。下面我们看下kernel的写法：

1: __kernel void vecadd(__global const float* A, __global const float* B, __global float* C)
2: {
3: int id = get_global_id(0);
4: C[id] = A[id] + B[id];
5: }

每个Kernel函数都必须以__kernel开始，而且必须返回void。每个输入参数都必须声明使用的内存类型。通过一些API，比如get_global_id之类的得到线程id。

内存对象地址空间标识符有以下几种：

__global – memory allocated from global address space

__constant – a special type of read-only memory

__local – memory shared by a work-group

__private – private per work-item memory

__read_only/__write_only – used for images

Kernel函数参数如果是内存对象，那么一定是__global,__local或者constant。

运行Kernel

   首先要设置线程索引空间的维数以及workgroup大小等。

   我们通过函数clEnqueueNDRangeKerne把Kernel放在一个队列里，但不保证它马上执行，OpenCL driver会管理队列，调度Kernel的执行。注意：每个线程执行的代码都是相同的，但是它们执行数据却是不同的。

   该函数把要执行的Kernel函数放在指定的命令队列中，globald大小（线程索引空间）必须指定，local大小（work group）可以指定，也可以为空。如果为空，则系统会自动根据硬件选择合适的大小。event_wait_list用来选定一些events，只有这些events执行完后，该kernel才可能被执行，也就是通过事件机制来实现不同kernel函数之间的同步。

   当Kernel函数执行完毕后，我们要把数据从device memory中拷贝到host memory中去。

释放资源：

    大多数的OpenCL资源都是指针，不使用的时候需要释放掉。当然，程序关闭的时候这些对象也会被自动释放掉。

    释放资源的函数是：clRelase{Resource} ，比如: clReleaseProgram(), clReleaseMemObject()等。

错误捕捉：

    如果OpenCL函数执行失败，会返回一个错误码，一般是个负值，返回0则表示执行成功。我们可以根据该错误码知道什么地方出错了，需要修改。错误码在cl.h中定义，下面是几个错误码的例子.

CL_DEVICE_NOT_FOUND -1

CL_DEVICE_NOT_AVAILABLE -2

CL_COMPILER_NOT_AVAILABLE -3

CL_MEM_OBJECT_ALLOCATION_FAILURE -4

…

下面是一个OpenCL机制的示意图

程序模型

    数据并行：work item和内存对象元素之间是一一映射关系；workgroup可以显示指定，也可以隐式指定。

    任务并行：kernel的执行独立于线程索引空间；用其他方法表示并行，比如把不同的任务放入队列，用设备指定的特殊的向量类型等等。

    同步：workgroup内work item之间的同步；命令队列中不同命令之间的同步。

完整代码如下：

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: using namespace std;
10: #define NWITEMS 262144
11:
12: #pragma comment (lib,"OpenCL.lib")
13:
14: //把文本文件读入一个string中
15: int convertToString(const char *filename, std::string& s)
16: {
17: size_t size;
18: char* str;
19:
20: std::fstream f(filename, (std::fstream::in | std::fstream::binary));
21:
22: if(f.is_open())
23: {
24: size_t fileSize;
25: f.seekg(0, std::fstream::end);
26: size = fileSize = (size_t)f.tellg();
27: f.seekg(0, std::fstream::beg);
28:
29: str = new char[size+1];
30: if(!str)
31: {
32: f.close();
33: return NULL;
34: }
35:
36: f.read(str, fileSize);
37: f.close();
38: str[size] = '\0';
39:
40: s = str;
41: delete[] str;
42: return 0;
43: }
44: printf("Error: Failed to open file %s\n", filename);
45: return 1;
46: }
47:
48: int main(int argc, char* argv[])
49: {
50: //在host内存中创建三个缓冲区
51: float *buf1 = 0;
52: float *buf2 = 0;
53: float *buf = 0;
54:
55: buf1 =(float *)malloc(NWITEMS * sizeof(float));
56: buf2 =(float *)malloc(NWITEMS * sizeof(float));
57: buf =(float *)malloc(NWITEMS * sizeof(float));
58:
59: //初始化buf1和buf2的内容
60: int i;
61: srand( (unsigned)time( NULL ) );
62: for(i = 0; i < NWITEMS; i++)
63: buf1[i] = rand()%65535;
64:
65: srand( (unsigned)time( NULL ) +1000);
66: for(i = 0; i < NWITEMS; i++)
67: buf2[i] = rand()%65535;
68:
69: for(i = 0; i < NWITEMS; i++)
70: buf[i] = buf1[i] + buf2[i];
71:
72: cl_uint status;
73: cl_platform_id platform;
74:
75: //创建平台对象
76: status = clGetPlatformIDs( 1, &platform, NULL );
77:
78: cl_device_id device;
79:
80: //创建GPU设备
81: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,
82: 1,
83: &device,
84: NULL);
85: //创建context
86: cl_context context = clCreateContext( NULL,
87: 1,
88: &device,
89: NULL, NULL, NULL);
90: //创建命令队列
91: cl_command_queue queue = clCreateCommandQueue( context,
92: device,
93: CL_QUEUE_PROFILING_ENABLE, NULL );
94: //创建三个OpenCL内存对象，并把buf1的内容通过隐式拷贝的方式
95: //拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2
96: cl_mem clbuf1 = clCreateBuffer(context,
97: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
98: NWITEMS*sizeof(cl_float),buf1,
99: NULL );
100:
101: cl_mem clbuf2 = clCreateBuffer(context,
102: CL_MEM_READ_ONLY ,
103: NWITEMS*sizeof(cl_float),NULL,
104: NULL );
105:
106: status = clEnqueueWriteBuffer(queue, clbuf2, 1,
107: 0, NWITEMS*sizeof(cl_float), buf2, 0, 0, 0);
108:
109: cl_mem buffer = clCreateBuffer( context,
110: CL_MEM_WRITE_ONLY,
111: NWITEMS * sizeof(cl_float),
112: NULL, NULL );
113:
114: const char * filename = "add.cl";
115: std::string sourceStr;
116: status = convertToString(filename, sourceStr);
117: const char * source = sourceStr.c_str();
118: size_t sourceSize[] = { strlen(source) };
119:
120: //创建程序对象
121: cl_program program = clCreateProgramWithSource(
122: context,
123: 1,
124: &source,
125: sourceSize,
126: NULL);
127: //编译程序对象
128: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );
129: if(status != 0)
130: {
131: printf("clBuild failed:%d\n", status);
132: char tbuf[0x10000];
133: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);
134: printf("\n%s\n", tbuf);
135: return -1;
136: }
137:
138: //创建Kernel对象
139: cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );
140: //设置Kernel参数
141: cl_int clnum = NWITEMS;
142: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);
143: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);
144: clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);
145:
146: //执行kernel
147: cl_event ev;
148: size_t global_work_size = NWITEMS;
149: clEnqueueNDRangeKernel( queue,
150: kernel,
151: 1,
152: NULL,
153: &global_work_size,
154: NULL, 0, NULL, &ev);
155: clFinish( queue );
156:
157: //数据拷回host内存
158: cl_float *ptr;
159: ptr = (cl_float *) clEnqueueMapBuffer( queue,
160: buffer,
161: CL_TRUE,
162: CL_MAP_READ,
163: 0,
164: NWITEMS * sizeof(cl_float),
165: 0, NULL, NULL, NULL );
166: //结果验证，和cpu计算的结果比较
167: if(!memcmp(buf, ptr, NWITEMS))
168: printf("Verify passed\n");
169: else printf("verify failed");
170:
171: if(buf)
172: free(buf);
173: if(buf1)
174: free(buf1);
175: if(buf2)
176: free(buf2);
177:
178: //删除OpenCL资源对象
179: clReleaseMemObject(clbuf1);
180: clReleaseMemObject(clbuf2);
181: clReleaseMemObject(buffer);
182: clReleaseProgram(program);
183: clReleaseCommandQueue(queue);
184: clReleaseContext(context);
185: return 0;
186: }
187:

也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下载完整版本。