在我们熟知的桌面平台，GPU得到了极为广泛的应用，小到各种电子游戏，大到高性能计算，多核心、高并行化的GPU成为我们日常娱乐和科学研究必不可少的“利器”。同样，在近些年兴起的移动平台，诸如智能手机、平板电脑等，也日渐重视GPU在其应用中的作用。近几年，随着并行化的发展，越来越多的手持设备硬件厂商重视对并行化标准的支持和应用。这里，需要支持OpenCL这一开发运算标准，该标准以异构平台为目标，与CUDA、Direct Compute主要面向PC平台不同，因而得到了众多厂商的支持，如下表：

（主要是高通的产品）

而在国外的一些研究机构和学者也对智能手机、平板电脑这样的移动平台进行了并行化的研究，比如三星手机研究院和诺基亚研究院近几年就发表了很多关于这方面的资料；美国莱斯大学的学者Guohui Wang等人就对物品移除算法和SIFT算法进行了智能手机上的并行化实现。

并行计算已经在移动平台具备硬件条件和变成标准的支持，而并行化又可以带来提升设备硬件利用效率，同时GPU的低主频特性又可以在一定程度上降低功耗，因此在智能手机等移动平台实现并行计算具有巨大的潜在价值，特别在当前手机续航时间不能满足用户要求的背景下，并行化的特性显得尤为重要。

下面就具体介绍一些实现流程和结果。

这次仅仅通过Sobel滤波这样的小程序来完成基于OpenCL实现的Android平台并行化。

首先，我们需要完成开发环境的搭建。由于目标是安卓平台，我们需要安装JAVA SDK、Android SDK、Eclipse以及ADT插件，这些工具的安装教程很多，这里就不再赘述了，主要介绍Cygwin与NDK的环境搭建。

第一步，从http://developer.android.com/tools/sdk/ndk/index.html查看和下载NDK工具的相关资料和安装包，我们在开发时使用的是NDK r8版本，后续版本使用基本类似；可以参照NDK的文档进行深入的学习和测试。

第二步，从http://www.cygwin.com 下载Cygwin工具。由于NDK完成的工作是允许开发人员使用本地代码(如C/C++)进行Android APP功能开发，而在开发的过程中大多涉及到GCC环境下的编译、运行，我们采用了Cygwin模拟Linux编译环境。我安装的时候为了方便就把所有的文件都安装了，体积不大，1G左右。

安装完成后，运行"Cygwin.bat"，可以通过以下几个方法检验安装是否成功。（这里参照了以前的一些资料）

（1）cygcheck –c cygwin（正常显示如图）

（2）make –v，返回make命令的版本信息

（3）gcc -v，返回gcc命令的版本信息

第三步，配置NDK的系统环境变量，为了避免编译时警告，可采用Linux风格的路径，如我的NDK安装路径为：“D:\android-ndk-r8-windows \android-ndk-r8”，在系统变量中名为“ndk”的变量路径为：“/cygdrive/ android-ndk-r8-windows/ android-ndk-r8”

另外，在Eclipse环境中可以安装CDT和Sequoyah插件方便Android工程对Native的开发（可以省略每次修改代码后都需要手动到代码目录进行ndk-build的步骤），可从http://www.eclipse.org/cdt/downloads.php下载CDT的离线安装包，然后再Eclipse中点击Help->Install New Software，点击Archive确定安装包所在位置，然后进行安装；Sequoyah可以直接在线安装，Location为：http://download.eclipse.org/sequoyah/updates/2.0/

如图： （安装时不要勾选“Group items by category”选项，否则会出现列表为空的情况）。然后在Window->preferences->Android->本机开发选项中添加NDK的安装路径。

其次，我会简要的介绍OpenCL在Android开发过程中的一些设置和代码。在Android平台实现并行化的过程中，我主要遵循下面的框图进行： 

主要思路就是利用JAVA中的JNI接口，结合Cygwin环境和NDK工具，将OpenCL实现的并行算法编译为可以被Android工程调用的libSobelFilter.so（lib***.so均可），然后在程序中调用该文件中的算法实现并行。

在Eclipse工程中jni文件夹下需要创建如下两个文件：Android.mk以及Application.mk，相当于该程序对应的makefile文件，前者定义了一系列规则来制定编译文件的目标和文件编译的顺序，后者定义了程序平台版本和编译器版本等内容。具体实现为，Android.mk文件：

[cpp] view plaincopy

1. LOCAL_PATH := $(call my-dir)
2. include $(CLEAR_VARS)
3. include ../../sdk/native/jni/OpenCV.mk #程序中利用了OpenCV加载图片
4. LOCAL_MODULE    := SobelFilter#生成的库文件名称
5. LOCAL_SRC_FILES := ImageSobelFilter.cpp aopencl.c#包含的文件
6. LOCAL_LDLIBS+=-L$(SYSROOT)/usr/lib -llog -ldl#库文件目录
7. include $(BUILD_SHARED_LIBRARY)

Application.mk文件

[html] view plaincopy

1. APP_CPPFLAGS := -frtti -fexceptions
2. APP_STL := stlport_static
3. APP_ABI := armeabi-v7a#程序的二进制接口
4. APP_PLATFORM := android-8#平台版本

然后将OpenCL头文件拷贝到jni文件夹下，供工程编译时调用：

接下来需要我们按照OpenCL的框架流程进行并行化的初始化和内核入队操作，主要包括：

1）获得平台clGetPlatformIDs；2）创建上下文clCreateContexFromType；3）通过上下文得到设备信息clGetContextInfo；4）为相应设备创建commandQueue；5）创建源程序，生成kernel；6）分配buffer空间，设置程序参数；7）执行kernel，clEnqueueNDRangeKernel；8）从buffer读回数据clEnqueueReadBuffer。

这里几个操作时OpenCL的固定流程，具体代码很多，请大家参看下我的源码，这里就不写了。

这里我要指出的是由于移动平台的特殊性，我们在程序中对几个宏变量进行了定义：

[cpp] view plaincopy

1. #define LIB_OPENCL "/system/vendor/lib/libOpenCL.so"//库文件
2. #define LIB_GLES_MALI "/system/vendor/lib/egl/libGLES_mali.so"//初始化
3. #define LIB_LLVM   "/system/vendor/lib/libllvm-a3xx.so"//构架编译器

上述几个Android平台需要的文件在不同版本的安卓系统中有不同的位置，上例

为Android 4.3版本的文件位置，在之前版本中文件多数位于"/system/lib/"文件夹下。

程序的核函数如下：

[cpp] view plaincopy

1. <span style=< span="">"font-size:18px;">__kernel void Sobel(__global char *array1, __global char *array2, __global int *array3)
2. {                                                                               size_t gidx = get_global_id(0);
3. size_t gidy = get_global_id(1); //此处的id为获得的设备序号
4. unsigned char a00, a01, a02;
5. unsigned char a10, a11, a12;
6. unsigned char a20, a21, a22;
7. int width=array3[0];                                                    int heigh=array3[1];                                                    int widthStep=array3[2];                        if(gidy>0&&gidy<heigh-1&&gidx>0&&gidx<width-1)                      {                                                                               a00 = array1[gidx-1+widthstep*(gidy-1)];< span="">//定义算子矩阵
8. a01 = array1[gidx+widthStep*(gidy-1)];
9. a02 = array1[gidx+1+widthStep*(gidy-1)];
10. a10 = array1[gidx-1+widthStep*gidy];
11. a11 = array1[gidx+widthStep*gidy];
12. a12 = array1[gidx+1+widthStep*gidy];                                    a20 = array1[gidx-1+widthStep*(gidy+1)];
13. a21 = array1[gidx+widthStep*(gidy+1)];
14. a22 = array1[gidx+1+widthStep*(gidy+1)];
15. float ux=a20+2*a21+a22-a00-2*a01-a02;   //定义卷积过程
16. float uy=a02+2*a12+a22-a00-2*a10-a20;
17. float u=sqrt(ux*ux + uy*uy);            //计算该点梯度
18. if(u>255) {
19. u=-1;
20. }
21. else if(u<20){
22. u=0;
23. }
24. array2[gidx+widthStep*gidy]=u;                                  }
25. }

此外，该程序使用了OpenCV的相关函数完成图像操作，因此需要在对应的安卓手机上安装OpenCV-Manager来完成对OpenCV函数的调用工作（可从http://opencv.org/下载相关资料和安装包）。同时，为了程序编译的方便，建议将程序文件放置到OpenCV-android-sdk的samples目录下，同时在Eclipse的项目属性Android选项中将＂…\OpenCV-2.4.8-android-sdk\sdk\java＂工程加入Android工程中，如图所示：

在完成运行环境的配置后，在Cygwin中完成该项目的编译和库文件生成工作，如下图所示：

可在项目工程目录下的libs/armeabi-v7a下查看生成的.so文件(libSobelFilter.so)

至此，Sobel滤波程序已经编译完成。下面介绍一些运行结果。

我们在DragonBoard开发板上进行程序测试。我们采用的常见的lena图像，改变图像的大小进行对比，这里我们采用了256*256,512*512,1024*1024等大小不同的图片进行测试，检验并行程序与串行的加速比。DragonBoard采用骁龙800处理器,同时，该开发板提供了丰富的板上资源，包括Snapdragon 800 Processor（4核2.15GHz，GPU为Adreno 330）、BT4.0、HDMI output、Dual SATA等，产自INTRINSYC公司(详情请参见：http://shop.intrinsyc.com/ ，目前已经有Snapdragon 805平台的开发板)。

以下的图表分别展示了程序的运行界面和加速比对比。

表中是一些测量得到的结果：

图片大小	并行时间（ms）	串行时间（ms）	加速比
256*256	4.95	7,67	1.55
512*512	27.62	44.27	1.58
1024*1024	88.86	138.66	1.57
2048*2048	154.41	241.65	1.56
4096*4096	247.65	468.28	1.71
5000*5000	693.92	1321.41	1.91

从上述结果可以看出，在上述实验平台上，随着图片大小的增大（数据处理更加复杂），并行化的加速比会更加明显。

在其他的智能手机，如小米2s（CPU：高通骁龙600，GPU：Adreno320）也做了类似的测试，也可以实现图片处理的加速。程序的工程文件如下：http://download.csdn.net/detail/u011723240/7485839

国创项目组成员：王帅  严章熙  周桐  （西安电子科技大学 电子工程学院）指导教师：朱虎明 副教授（西电智能所）

致谢：感谢西安电子科技大学国家大学生创新实验计划的支持（编号201310701017）、感谢“AMD-西电多核异构高性能实验室”提供部分设备支持