C语言与OpenCL的编程示例比较
C语言与OpenCL的编程示例比较
OpenCL支持数据并行，任务并行编程，同时支持两种模式的混合。对于同步 OpenCL支持同一工作组内工作项的同步和命令队列中处于同一个上下文中的 命令的同步。
在本文中以图像旋转的实例，具体介绍OpenCL编程的步骤。 首先给出实现流程，然后给出实现图像旋转的C循环实现和OpenCL C kernel实现。
4.2 图像旋转
4.2.1 图像旋转原理
图像旋转是指把定义的图像绕某一点以逆时针或顺时针方向旋转一定的角度， 通常是指绕图像的中心以逆时针方向旋转。假设图像的左上角为(l, t), 右下角为(r, b)，则图像上任意点(x, y) 绕其中心(xcenter, ycenter)逆时针旋转θ角度后， 新的坐标位置(x',y')的计算公式为：
x′ = (x - xcenter) cosθ - (y － ycenter) sinθ + xcenter,
y′ = (x - xcenter) sinθ + (y － ycenter) cosθ + ycenter.
C代码：
void rotate(
unsigned char* inbuf,
unsigned char* outbuf,
int w, int h,
float sinTheta,
float cosTheta)
{
int i, j;
int xc = w/2;
int yc = h/2;
for(i = 0; i < h; i++)
{
for(j=0; j< w; j++)
{
int xpos = (j-xc)*cosTheta - (i - yc) * sinTheta + xc;
int ypos = (j-xc)*sinTheta + (i - yc) * cosTheta + yc;
if(xpos>=0&&ypos>=0&&xpos<w&&ypos<h)
outbuf[ypos*w + xpos] = inbuf[i*w+j];
}
}
}
OpenCL C kernel代码：
#pragma OPENCL EXTENSION cl_amd_printf : enable
__kernel void image_rotate(
__global uchar * src_data,
__global uchar * dest_data, //Data in global memory
int W, int H, //Image Dimensions
float sinTheta, float cosTheta ) //Rotation Parameters
{
const int ix = get_global_id(0);
const int iy = get_global_id(1);
int xc = W/2;
int yc = H/2;
int xpos = ( ix-xc)*cosTheta - (iy-yc)*sinTheta+xc;
int ypos = (ix-xc)*sinTheta + ( iy-yc)*cosTheta+yc;
if ((xpos>=0) && (xpos< W) && (ypos>=0) && (ypos< H))
dest_data[ypos*W+xpos]= src_data[iy*W+ix];
}
正如上面代码中所给出的那样，在C代码中需要两重循环来计算横纵坐标上新的 坐标位置。其实，在图像旋转的算法中每个点的计算可以独立进行，与其它点的 坐标位置没有关系，所以并行处理较为方便。OpenCL C kernel代码中用了并行 处理。
上面的代码在Intel的OpenCL平台上进行了测试，处理器为双核处理器，图像大小 为4288*3216，如果用循环的方式运行时间稳定在0.256s左右，而如果用OpenCL C kernel并行的方式，运行时间稳定在0.132秒左右。GPU的测试在NVIDIA的GeForce G105M显卡 上进行，运行时间稳定在0.0810s左右。从循环的方式，双核CPU并行以及GPU并行计算 已经可以看出，OpenCL编程的确能大大提高执行效率。
通过对OpenCL编程的分析和实验可以得出，用OpenCL编写的应用具有很好的移 植性，能在不同的设备上运行。OpenCL C kernel一般用并行的方式处理，所以能极大地提高程序的运行效率。
内存模型
一般而言，不同的平台之间有不同的存储系统。例如，CPU有高速缓存而GPU就没有。 为了程序的可移植性，OpenCL定义了抽象的内存模型，程序实现的时候只需关注抽 象的内存模型，具体向硬件上的映射由驱动来完成。内存空间的定义及与硬件的映 射大致如图所示。
内存空间在程序内部可以用关键字的方式指定，不同的定义与数据存在的位置 相关，主要有如下几个基本概念[ 2 ]：
- 全局内存：所有工作组中的所有工作项都可以对其进行读写。工作项可以 读写此中内存对象的任意元素。对全局内存的读写可能会被缓存，这取决于设备 的能力。
- 不变内存：全局内存中的一块区域，在内核的执行过程中保持不变。 宿主机负责对此中内存对象的分配和初始化。
- 局部内存：隶属于一个工作组的内存区域。它可以用来分配一些变量， 这些变量由此工作组中的所有工作项共享。在OpenCL设备上，可能会将 其实现成一块专有的内存区域，也可能将其映射到全局内存中。
- 私有内存：隶属于一个工作项的内存区域。 一个工作项的私有内存中所定义的变量对另外一个工作项来说是不可见的。
参考链接：
https://www.cnblogs.com/wangshide/archive/2012/01/07/2315830.html
【1】 Aaftab Munshi. The OpenCL Specification Version1.1 Document Revision:44[M]. Khronos OpenCL Working Group. 2011.6.1.
【2】Aaftab Munshi. 倪庆亮译. OpenCL规范 Version1.0 Document Revision:48[M]. Khronos OpenCL Working Group. 2009.10.6.
【3】Aaftab Munshi, Benedict R. Gaster, Timothy G. Mattson, James Fung, Dan Ginsburg. OpenCL Programming Guide [M]. Addison-Wesley Professional. 2011.7.23.
【4】Benedict Gaster, Lee Howes, David R. Kaeli and Perhaad Mistry. Heterogeneous Computing with OpenCL[M]. Morgan Kaufmann, 1 edition. 2011.8.31.
【5】Slo-Li Chu, Chih-Chieh Hsiao. OpenCL: Make Ubiquitous Supercomputing Possible[J]. IEEE International Conference on High Performance Computing and Communications. 2010 12th 556-561.
【6】John E. Stone, David Gohara, Guochun Shi. OpenCL: A parallel programming standard for heterogeneous computing systems[J]. Copublished by the IEEE CS and the AIP. 2010.5/6 66-72.
【7】Kyle Spafford, Jeremy Meredith, Jeffrey Vetter. Maestro:Data Orchestration and Tuning for OpenCL Devices[J]. P. D'Ambra,M.Guarracino, and D.Talia (Eds.):Euro-Par 2010,Part II,LNCS6272, pp. 275–286, 2010. \copyright Springer-Verlag Berlin Heidelberg 2010.
