Android平臺利用OpenCL框架實現並行開發初試
款式 | CPU型號 | GPU型號 | OpenCL支援 |
三星GalaxyS5 | 高通驍龍801(4核) | Adreno330 | 是 |
Iphone5S | 蘋果A7(2核) | Imagination PowerVR G6430 | 是 |
小米3 | 高通驍龍800(4核) | Adreno330 | 是 |
魅族M3 | 三星5410(8核) | Imagination SGX544 | 是 |
而在國外的一些研究機構和學者也對智慧手機、平板電腦這樣的移動平臺進行了並行化的研究,比如三星手機研究院和諾基亞研究院近幾年就發表了很多關於這方面的資料;美國萊斯大學的學者Guohui Wang等人就對物品移除演算法和SIFT演算法進行了智慧手機上的並行化實現。
平行計算已經在移動平臺具備硬體條件和變成標準的支援,而並行化又可以帶來提升裝置硬體利用效率,同時GPU的低主頻特性又可以在一定程度上降低功耗,因此在智慧手機等移動平臺實現平行計算具有巨大的潛在價值,特別在當前手機續航時間不能滿足使用者要求的背景下,並行化的特性顯得尤為重要。
下面就具體介紹一些實現流程和結果。
這次僅僅通過Sobel濾波這樣的小程式來完成基於OpenCL實現的Android平臺並行化。
首先,我們需要完成開發環境的搭建。由於目標是安卓平臺,我們需要安裝JAVA SDK、Android SDK、Eclipse以及ADT外掛,這些工具的安裝教程很多,這裡就不再贅述了,主要介紹Cygwin與NDK的環境搭建。
第二步,從http://www.cygwin.com 下載Cygwin工具。由於NDK完成的工作是允許開發人員使用原生代碼(如C/C++)進行Android
APP功能開發,而在開發的過程中大多涉及到GCC環境下的編譯、執行,我們採用了Cygwin
安裝完成後,執行"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確定安裝包所在位置,然後進行安裝;
如圖:(安裝時不要勾選“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檔案:
LOCAL_PATH := $(call my-dir)
include $(CLEAR_VARS)
include ../../sdk/native/jni/OpenCV.mk #程式中利用了OpenCV載入圖片
LOCAL_MODULE := SobelFilter#生成的庫檔名稱
LOCAL_SRC_FILES := ImageSobelFilter.cpp aopencl.c#包含的檔案
LOCAL_LDLIBS+=-L$(SYSROOT)/usr/lib -llog -ldl#庫檔案目錄
include $(BUILD_SHARED_LIBRARY)
Application.mk檔案
APP_CPPFLAGS := -frtti -fexceptions
APP_STL := stlport_static
APP_ABI := armeabi-v7a#程式的二進位制介面
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的固定流程,具體程式碼很多,請大家參看下我的原始碼,這裡就不寫了。
這裡我要指出的是由於移動平臺的特殊性,我們在程式中對幾個巨集變數進行了定義:
#define LIB_OPENCL "/system/vendor/lib/libOpenCL.so"//庫檔案
#define LIB_GLES_MALI "/system/vendor/lib/egl/libGLES_mali.so"//初始化
#define LIB_LLVM "/system/vendor/lib/libllvm-a3xx.so"//構架編譯器
上述幾個Android平臺需要的檔案在不同版本的安卓系統中有不同的位置,上例
為Android 4.3版本的檔案位置,在之前版本中檔案多數位於"/system/lib/"資料夾下。
程式的核函式如下:
<span style="font-size:18px;">__kernel void Sobel(__global char *array1, __global char *array2, __global int *array3)
{ size_t gidx = get_global_id(0);
size_t gidy = get_global_id(1); //此處的id為獲得的裝置序號
unsigned char a00, a01, a02;
unsigned char a10, a11, a12;
unsigned char a20, a21, a22;
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)];//定義運算元矩陣
a01 = array1[gidx+widthStep*(gidy-1)];
a02 = array1[gidx+1+widthStep*(gidy-1)];
a10 = array1[gidx-1+widthStep*gidy];
a11 = array1[gidx+widthStep*gidy];
a12 = array1[gidx+1+widthStep*gidy]; a20 = array1[gidx-1+widthStep*(gidy+1)];
a21 = array1[gidx+widthStep*(gidy+1)];
a22 = array1[gidx+1+widthStep*(gidy+1)];
float ux=a20+2*a21+a22-a00-2*a01-a02; //定義卷積過程
float uy=a02+2*a12+a22-a00-2*a10-a20;
float u=sqrt(ux*ux + uy*uy); //計算該點梯度
if(u>255) {
u=-1;
}
else if(u<20){
u=0;
}
array2[gidx+widthStep*gidy]=u; }
}
此外,該程式使用了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 |
從上述結果可以看出,在上述實驗平臺上,隨著圖片大小的增大(資料處理更加複雜),並行化的加速比會更加明顯。
國創專案組成員:王帥 嚴章熙 周桐 (西安電子科技大學 電子工程學院)指導教師:朱虎明 副教授(西電智慧所)
致謝:感謝西安電子科技大學國家大學生創新實驗計劃的支援(編號201310701017)、感謝“AMD-西電多核異構高效能實驗室”提供部分裝置支援