在Android上使用OpenCL調用GPU加速


在Android上還實現了很多種並行化的算法,比如SHA-1、HDR、K-means、NL-means、SRAD等等,會在近期整理好之后開源的。
原文發表在了 異構開發技術社區
整理成教程是隊友做的,十分感謝~
原博文地址
隊友的博客

 

下面是干貨:

Android平台利用OpenCL框架實現並行開發初試

在我們熟知的桌面平台,GPU得到了極為廣泛的應用,小到各種電子游戲,大到高性能計算,多核心、高並行化的GPU成為我們日常娛樂和科學研究必不可少的“利器”。同樣,在近些年興起的移動平台,諸如智能手機、平板電腦等,也日漸重視GPU在其應用中的作用。近幾年,隨着並行化的發展,越來越多的手持設備硬件廠商重視對並行化標准的支持和應用。這里,需要支持OpenCL這一開發運算標准,該標准以異構平台為目標,與CUDA、Direct Compute主要面向PC平台不同,因而得到了眾多廠商的支持,如下表:

常見智能手機的硬件信息款式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://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文件:
 
Application.mk文件:
 
然后將OpenCL頭文件拷貝到jni文件夾下,供工程編譯時調用:
 
接下來需要我們按照OpenCL的框架流程進行並行化的初始化和內核入隊操作,主要包括:
1)獲得平台clGetPlatformIDs;2)創建上下文clCreateContexFromType;3)通過上下文得到設備信息clGetContextInfo;4)為相應設備創建commandQueue;5)創建源程序,生成kernel;6)分配buffer空間,設置程序參數;7)執行kernel,clEnqueueNDRangeKernel;8)從buffer讀回數據clEnqueueReadBuffer。
這里幾個操作時OpenCL的固定流程,具體代碼很多,請大家參看下我的源碼,這里就不寫了。
這里我要指出的是由於移動平台的特殊性,我們在程序中對幾個宏變量進行了定義:
 
上述幾個Android平台需要的文件在不同版本的安卓系統中有不同的位置,上例
為Android 4.3版本的文件位置,在之前版本中文件多數位於”/system/lib/”文件夾下。
程序的核函數如下:
 
此外,該程序使用了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)也做了類似的測試,也可以實現圖片處理的加速


免責聲明!

本站轉載的文章為個人學習借鑒使用,本站對版權不負任何法律責任。如果侵犯了您的隱私權益,請聯系本站郵箱yoyou2525@163.com刪除。



 
粵ICP備18138465號   © 2018-2025 CODEPRJ.COM