Apple's OpenCL——以一个简单的实例介绍各个基本概念(转)
http://www.cocoachina.com/bbs/read.php?tid-31352.html
首先，感谢各位，感谢CocoaChina对OpenCL的关注。看到首页上一个大大的OpenCL图标，在下着实激动不已，呵呵。
另外，再次感谢Apple能够搞出OpenCL来方便地利用GPU等加速处理器来解决高兴能计算领域的问题。
在《概述》章节中，我把OpenCL中遇到的一些基本概念以提纲的形式罗列了出来，目的在于给各位以一个框架性的介绍，然后各位可以在网上搜索到相关的更具体的介绍。
就目前对于Mac而言，要获悉，或者说要对OpenCL的工作机制有更清晰的认识，那么有必要先对GPU在高性能计算上的一个处理特征做一些了解。
我这里提供一个链接给大家提供参考。
http://vga.zol.com.cn/170/1709097.html
这个链接中还有其它文章对GPU从不同角度的介绍。就目前而言，CPU仅有的那些ALU以及用于SIMD的计算单元不足以支撑其高性能计算（HPC）的重任。因此，我们可以看到历年被评为世界级速度最快的超级计算机一般有多个处理器，多个图形处理器等组成，功耗是高得一塌糊涂。目前，还有很多传统的做法是通过分布式计算来提高计算性能。由此可见，HPC就目前而言需求还是非常广的。
为什么GPU比CPU更能胜任高性能计算？因为GPU的架构以及设计理念与CPU不同。GPU生来就是用于处理视频解码、图像处理以及2D、3D图形处理的。而这些应用都有一个特点——比如，举个简单的例子，对一个图像做半透明效果，即图像颜色与背景色做混合计算。由于在该计算中，每对像素的乘加都是相互独立的，因此对于同一个乘加操作能作用于整幅图片，这样理所当然地就能利用SIMD特性去处理——即，单条指令作用于多个数据。如果熟悉x86架构中SIMD特性的朋友应该会注意到SSE也好，MMX也罢，这些操作都需要将标量数据打包，放进一个SSE或MMX寄存器中，然后对这些64位、128位寄存器做计算操作后，最后还要将数据拆包，放到32位寄存器或者是存储器中。而GPU就没有这种过程。它是天然地支持SIMD。
各位在上述链接中第三页能看到nVidia最新GF100的架构，称为Fermi（费米）。在第四页能看到TPC的结构图。
这里我们将引入两个概念，一个是“Device”（设备），一块GPU可以看作为一台设备，当然一块Intel Core CPU也能被看作为一台OpenCL计算设备；
另一个概念是“Compute Unit”（计算单元）。计算单元是处理器中用于解析指令并进行执行的核心单元。比如x86的CPU中，一个计算单元，在物理上就是一个Core（核心）；而在nVidia的GPU中，一个计算单元，在物理上就是一个SM（Stream Multiprocessor，即流多处理器）。我们在第4页能看到在G80架构中，一个TPC含有两个SM，即两个计算单元；而GT200架构中，则有3个。我的Mac Mini的GPU用的是GeForce 9400，属于G80架构，由于只有一个TPC，因此仅有两个计算单元……:(
一般来说，一个计算单元可能还包含多个线程。比如，x86中如果处理器拥有超线程技术，那么一个核还可以分成两个逻辑处理器。而对于G80架构的GPU而言，一个SM中还含有8个SP（Stream Processor，即流标量处理器）。而一个计算单元中的线程有一个共享存储器，x86的具有超线程的一个核中，两个逻辑核共享L1 Data Cache以及L1 Instruction Cache；在G80中每个SM的所有SP共享一块shared memory。并且每个线程有私有存储器，比如x86中，每个逻辑核有独立的寄存器文件；在G80中，每个SP有自己的Local Memroy。
以上我对计算单元的一个整体介绍结合了CPU与GPU本身的特性进行。由于nVidia的Tesla基于CUDA的架构能直接映射到OpenCL的各个元素概念，因此我们以后将以此架构进行介绍。
下面将正式进入我们这一讲的主题，下面的代码来自于苹果开发者网站中，Apple自己的一个OpenCL Demo程序。
单单用来高清和游戏的话PCIe基本差不多了，但是要对付HPC的话……必须离内存更加近。
>CPU仅有的那些ALU以及用于SIMD的计算单元不足以支撑其高性能计算（HPC）的重任。
单看裸跑马力CPU当然比不过GPU。但是现在CPU代码移植到GPU上速度反而更慢的应用有一堆。没有人力去移植的又有一堆。这样那些top500，全都堆CPU，但不是全都堆GPU。
我们用我在概述中发的OpenCL_Query程序中会发现，GPU设备的CL_DEVICE_MAX_WORK_GROUP_SIZE有上百个，比如我的Mac Mini中的Geforce9400就有512个，而CPU则只有1个。为什么双核CPU却只有一个work-item？
首先，在OpenCL中，每个work都是独立的，因此它不具备CPU多核协同处理的概念，你可以将不同的作业交给不同的CPU核去执行，这个模式被称为任务并行模式（Task Parallel）。而对于多个工作项执行同一份内核函数（Kernel Function）的模式被称为数据并行执行模式（Data Parallel）。OpenCL兼容的GPU具有原生的数据并行执行引擎，它的执行架构是给每个工作项执行同一个内核实例，每个工作项彼此独立，并对应不同的数据元素，因此也被称为单指令多数据执行模式（Single Instruction Multiple Data）。而CPU不具备这个执行特性，尽管它可能有向量类型的执行单元，但是OpenCL不会利用CPU的向量执行单元而对它做SIMD操作。主要原因是处理起来非常麻烦，有时，若内核函数比较复杂的话几乎就是无法实现的。比如几个工作项之间需要共享一些数据并需要同步等等。
因此，我们可以看到目前CPU的工作组大小、每个工作组的每个维度上的工作项的个数都是1。
下面，我们将以一个简单的例子来说明如何创建一个OpenCL应用，并且运行。
OpenCL应用分两个部分：一个是主机端应用；一个是OpenCL内核。
为了能在后面更方便地为各位做出解释，我将以带行号的方式将代码展示出来。
1 /*
2 * hello.c
3 * OpenCL_init
4 *
5 * Created by Zenny Chen on 9/1/10.
6 * Copyright 2010 GreenGames Studio. All rights reserved.
7 *
8 */
9
10 #include <fcntl.h>
11 #include <stdio.h>
12 #include <stdlib.h>
13 #include <string.h>
14 #include <math.h>
15 #include <unistd.h>
16 #include <sys/types.h>
17 #include <sys/stat.h>
18 #include <OpenCL/opencl.h>
19
20 ////////////////////////////////////////////////////////////////////////////////
21
22 // Use a static data size for simplicity
23 //
24 #define DATA_SIZE (1024)
25
26 ////////////////////////////////////////////////////////////////////////////////
27
28 // Simple compute kernel which computes the square of an input array
29 //
30 const char *KernelSource = "\n" \
31 "__kernel square( \n" \
32 " __global float* input, \n" \
33 " __global float* output, \n" \
34 " const unsigned int count) \n" \
35 "{ \n" \
36 " int i = get_global_id(0); \n" \
37 " if(i < count) \n" \
38 " output[i] = input[i] * input[i]; \n" \
39 "} \n" \
40 "\n";
41
42 ////////////////////////////////////////////////////////////////////////////////
43
44 int main(int argc, char** argv)
45 {
46 int err; // error code returned from api calls
47
48 float data[DATA_SIZE]; // original data set given to device
49 float results[DATA_SIZE]; // results returned from device
50 unsigned int correct; // number of correct results returned
51
52 size_t global; // global domain size for our calculation
53 size_t local; // local domain size for our calculation
54
55 cl_device_id device_id; // compute device id
56 cl_context context; // compute context
57 cl_command_queue commands; // compute command queue
58 cl_program program; // compute program
59 cl_kernel kernel; // compute kernel
60
61 cl_mem input; // device memory used for the input array
62 cl_mem output; // device memory used for the output array
63
64
65 // Fill our data set with random float values
66 //
67 int i = 0;
68 unsigned int count = DATA_SIZE;
69 for(i = 0; i < count; i++)
70 data[i] = rand() / (float)RAND_MAX;
71
72 // Connect to a compute device
73 //
74 int gpu = 1;
75 err = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);
76 if (err != CL_SUCCESS)
77 {
78 printf("Error: Failed to create a device group!\n");
79 return EXIT_FAILURE;
80 }
81
82 // Create a compute context
83 //
84 context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
85 if (!context)
86 {
87 printf("Error: Failed to create a compute context!\n");
88 return EXIT_FAILURE;
89 }
90
91 // Create a command commands
92 //
93 commands = clCreateCommandQueue(context, device_id, 0, &err);
94 if (!commands)
95 {
96 printf("Error: Failed to create a command commands!\n");
97 return EXIT_FAILURE;
98 }
99
100 // Create the compute program from the source buffer
101 //
102 program = clCreateProgramWithSource(context, 1, (const char **) & KernelSource, NULL, &err);
103 if (!program)
104 {
105 printf("Error: Failed to create compute program!\n");
106 return EXIT_FAILURE;
107 }
108
109 // Build the program executable
110 //
111 err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
112 if (err != CL_SUCCESS)
113 {
114 size_t len;
115 char buffer[2048];
116
117 printf("Error: Failed to build program executable!\n");
118 clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
119 printf("%s\n", buffer);
120 exit(1);
121 }
122
123 // Create the compute kernel in the program we wish to run
124 //
125 kernel = clCreateKernel(program, "square", &err);
126 if (!kernel || err != CL_SUCCESS)
127 {
128 printf("Error: Failed to create compute kernel!\n");
129 exit(1);
130 }
131
132 // Create the input and output arrays in device memory for our calculation
133 //
134 input = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * count, NULL, NULL);
135 output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(float) * count, NULL, NULL);
136 if (!input || !output)
137 {
138 printf("Error: Failed to allocate device memory!\n");
139 exit(1);
140 }
141
142 // Write our data set into the input array in device memory
143 //
144 err = clEnqueueWriteBuffer(commands, input, CL_TRUE, 0, sizeof(float) * count, data, 0, NULL, NULL);
145 if (err != CL_SUCCESS)
146 {
147 printf("Error: Failed to write to source array!\n");
148 exit(1);
149 }
150
151 // Set the arguments to our compute kernel
152 //
153 err = 0;
154 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
155 err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
156 err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);
157 if (err != CL_SUCCESS)
158 {
159 printf("Error: Failed to set kernel arguments! %d\n", err);
160 exit(1);
161 }
162
163 // Get the maximum work group size for executing the kernel on the device
164 //
165 err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(local), &local, NULL);
166 if (err != CL_SUCCESS)
167 {
168 printf("Error: Failed to retrieve kernel work group info! %d\n", err);
169 exit(1);
170 }
171
172 // Execute the kernel over the entire range of our 1d input data set
173 // using the maximum number of work group items for this device
174 //
175 global = count;
176 err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
177 if (err)
178 {
179 printf("Error: Failed to execute kernel!\n");
180 return EXIT_FAILURE;
181 }
182
183 // Wait for the command commands to get serviced before reading back results
184 //
185 clFinish(commands);
186
187 // Read back the results from the device to verify the output
188 //
189 err = clEnqueueReadBuffer( commands, output, CL_TRUE, 0, sizeof(float) * count, results, 0, NULL, NULL );
190 if (err != CL_SUCCESS)
191 {
192 printf("Error: Failed to read output array! %d\n", err);
193 exit(1);
194 }
195
196 // Validate our results
197 //
198 correct = 0;
199 for(i = 0; i < count; i++)
200 {
201 if(results[i] == data[i] * data[i])
202 correct++;
203 }
204
205 // Print a brief summary detailing the results
206 //
207 printf("Computed '%d/%d' correct values!\n", correct, count);
208
209 // Shutdown and cleanup
210 //
211 clReleaseMemObject(input);
212 clReleaseMemObject(output);
213 clReleaseProgram(program);
214 clReleaseKernel(kernel);
215 clReleaseCommandQueue(commands);
216 clReleaseContext(context);
217
218 return 0;
219 }
大家请先看第30到40行。这段代码就是给OpenCL设备执行用的。这个我们一般称之为内核函数。这里，它是以字符串的形式给出的。OpenCL驱动将会把这部分代码编译为一个OpenCL中间语言，然后根据不同的设备（CPU、GPU、DSP等），构造成最终的机器码，提供给目标设备执行。
这个内核函数所完成的功能非常简单，它将我们所输入的一个一维数组做开方运算，将每个元素的结果送到相应的输出缓存。输入和输出缓存均在主机端应用程序上构建并初始化。
从第44行开始，main函数里就是执行主机端的代码。
下面将为大家逐一介绍每步过程。
一、获得设备ID：
我们调用下列函数来获取一个或多个设备ID：
cl_int clGetDeviceIDs (cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries,
cl_device_id * devices,
cl_uint *num_devices
)
这里的第一个参数，platform是通过调用clGetPlatformIDs函数获得的平台ID。如果这个参数为空，那么行为是实现定义的。对于Apple的OpenCL驱动而言，这个参数传空，那么将默认使用Apple的OpenCL驱动；否则的话可以使用其它第三方的OpenCL驱动。目前，nVidia、AMD都有针对自己GPU所推出的OpenCL的驱动。因此在一台主机上有两个，甚至更多的OpenCL驱动都完全有可能——比如你的一台Mac上既有AMD的显卡，又有nVidia的显卡……
我们在第75行可以看到，这里其实传递的是空指针。
这里的第二个参数，device_type是个枚举类型。它有这么几个值：
CL_DEVICE_TYPE_CPU——表示将用于OpenCL的内核函数在CPU上运行；
CL_DEVICE_TYPE_GPU——表示将用于OpenCL的内核函数在一台GPU设备上运行；
CL_DEVICE_TYPE_ACCELERATOR——OpenCL专用加速器，比如IBM的CELL处理器
CL_DEVICE_TYPE_DEFAULT——系统默认的OpenCL设备
CL_DEVICE_TYPE_ALL——所有当前系统可用的OpenCL设备。
第三个参数，num_entries表示第四个参数devices所指数组的最大元素个数。
第四个参数，devices，指向一个设备ID数组的首地址。该函数所返回的设备ID都会依次放入该指针所指向的数组中。
第五个参数，num_devices，这是一个输出参数。如果这个参数不指向空，那么在调用完这个函数后，会把实际放入到设备ID数组中的元素个数写回到这个指针所指向的变量中。
我们参考75行，这里我们就指定了一个GPU设备，因此device_id其实是一个变量，而不是数组。因为最后写回的实际设备ID个数不会超过我们所指定的1。
当我们调用完这个函数后，如果返回值为成功，那么我们后面就可以用设备ID来创建上下文了。
二、获得上下文：
一个OpenCL上下文可以用一个或多个设备进行创建。它被OpenCL运行时用于管理对象，诸如命令队列、存储器、程序以及内核对象，并且管理在上下文中所指定的一个或多个设备上的执行内核。
下面给出函数原型：
cl_context clCreateContext (const cl_context_properties * properties, cl_uint num_devices,
const cl_device_id * devices,
void (*pfn_notify)(const char*errinfo,
const void *private_info, size_t cb, void *user_data),
void user_data,
cl_int *errcode_ret
)
下面介绍各个参数：
properties：指定一列上下文属性名和其相应的值。每个属性名后面紧跟其相应的值。列表以0结尾。在当前的OpenCL1.0中，支持的属性名只有CL_CONTEXT_PLATFORM一个，而其相对应的属性值则是一个cl_platform_id。
num_devices：指定传入设备数组的长度。这个值表征了第三个参数devices的可用最大长度。
devices：指向通过上面调用的clGetDeviceIDs所获得的设备列表的首地址。
pfn_notify：它是一个应用程序可注册的回调函数。该回调函数将被OpenCL实现所使用，用于在该上下文中错误发生时报告错误信息。回调函数可以被OpenCL实现异步地调用。应用程序应当确保该回调函数是线程安全的。
该回调函数的参数介绍：
｛
errinfo——是一个指向错误字符串的指针
private_info和cb表示由OpenCL实现所返回的一个指向二进制数据的指针。它可用于记录对调试有帮助的额外信息。
user_data——指向一块用户所提供的数据
｝
user_data——当pfn_notify被调用时，它将作为pfn_notify的最后一个实参。
下面将结合clCreateContext的使用，给出一份代码：
cl_platform_id platform_id;// added by zenny_chen
cl_device_id device_id; // compute device id
cl_context context;// compute context
// Create a platform
err = clGetPlatformIDs(1, &platform_id, NULL);
if (err != CL_SUCCESS)
{ printf("Error: Failed to create a platform!\n");
return EXIT_FAILURE;
}
// Connect to a compute device //
err = clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
if (err != CL_SUCCESS)
{ printf("Error: Failed to create a device group!\n");
return EXIT_FAILURE;
}
// Create a compute context //
context = clCreateContext((cl_context_properties[3]){(cl_context_properties)CL_CONTEXT_PLATFORM, (cl_context_properties)platform_id, 0},
1, &device_id, NULL, NULL, &err);
if (!context)
{
printf("Error: Failed to create a compute context!\n");
return EXIT_FAILURE;
}
三、创建一个命令队列：
OpenCL命令队列用于将工作提交给设备。它们安排一个设备上内核的执行，并操作存储器对象。
OpenCL按照命令的入队顺序执行命令。当然，命令也可以以无序模式进行执行。这个将在后面针对命令队列做详细介绍时再深入讲解。
创建一个OpenCL命令队列使用以下函数：
cl_command_queue clCreateCommandQueue (cl_context context, cl_device_id device,
cl_command_queue_properties properties,
cl_int * errcode_ret)
下面介绍参数：
context——这个就是我们在上面创建好的一个上下文，将它传入
device——我们指定好的一个device
properties——当前，OpenCL支持两种特性，一个是CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE，这个特性指示我们将要创建的命令队列将以无序的方式执行命令；另一个是CL_QUEUE_PROFILING_ENABLE，这个特性表示对将要创建的命令队列开启剖析功能，剖析功能一般用于性能测试。
errcode_ret——输出错误码
这个函数将返回一个命令队列句柄。
我们在示例代码的第93行看到，properties的实参是0，说明我们既不开启无序执行特性也不开启剖析特性。
四、用源代码缓存创建程序：
一个OpenCL程序是由一组内核构成。一个内核是源代码中用__kernel限定符修饰的函数。OpenCL程序也可包含辅助函数和常量数据。它们被内核函数所使用。
一个程序对象包含了以下信息：
1、一个相关联的上下文
2、一个程序源代码或二进制代码
3、最近成功被构建的可执行程序。该可执行程序是根据前面所创建好的设备列表而被构建。
4、当前所连接的内核对象的个数
我们可以通过调用clCreateProgramWithSource或clCreateProgramWithBinary来创建程序对象。我们这里将介绍clCreateProgramWithSource函数：
cl_program clCreateProgramWithSource (cl_context context, cl_uint count, const char **strings, const size_t *lengths, cl_int *errcode_ret)
参数介绍：
context——程序相关联的上下文
count——用于指定下面所传入的源代码字符串的个数
strings——指向源代码字符串数组（即指向一个const char*的数组的首地址），该数组中存放的是源代码字符串
lengths——为每个源代码字符串提供相应的长度。当然，如果我们上面每个字符串能够保证是以'\0'结尾，那么我们可把这个参数置空
errcode_ret——返回错误码
我们可以在3楼代码中的102行看到这个函数的具体使用。这里，传入的是仅有一个字符串元素，&KernelSource。KernelSource里面的内容就是要在OpenCL设备上运行的内核函数的源代码。
五、构建OpenCL程序：
我们使用clCreateProgramWithSource创建完的是一个包含源代码的程序对象。我们要运行设备上的OpenCL程序必须对它进行编译、连接。而对于我们用clCreateProgramWithBinary创建的程序对象，则需要进行连接。最后构建成一个完整的可加载的执行程序。
我们通过clBuildProgram这个函数来构建OpenCL程序：
cl_int clBuildProgram (cl_program program, cl_uint num_devices, const cl_device_id *device_list, const char *options, void (*pfn_notify)(cl_program, void *user_data), void *user_data)
下面介绍各个参数：
program——我们上面所创建的程序对象
num_devices——指定设备的个数；这个参数主要是指明了第三个参数device_list的长度
device_list——指向与程序对象相关联的设备列表
options——指向一个字符串，用于描述构建选项（这个就像是我们用控制台去编译一个源文件时用的编译选项一样）
pfn_notify——是一个指向通知例程的函数指针。这个参数与第二步，创建上下文时的那个参数功能一样。
