自從首部拍照手機問(wèn)世后,手機就成為視覺(jué)導向設備,今天用戶(hù)產(chǎn)生的內容已成為從網(wǎng)絡(luò )利用率到創(chuàng )建應用程序的主要推手,然而,開(kāi)發(fā)人員們最近開(kāi)始對圖像數據進(jìn)行高級處理,以創(chuàng )建一系列新的用戶(hù)體驗,這些體驗從智能視覺(jué)系統、增強現實(shí)應用程序到虛擬制作等等。不過(guò)這些創(chuàng )新技術(shù)的爆發(fā)受到手機數據處理量的限——這當中存在很大挑戰。 計算攝影:移動(dòng)計算的主要應用之一 成像一直以來(lái)是CPU的主任務(wù),對于靜止圖像和小格式視頻來(lái)說(shuō)CPU處理得很好,但隨著(zhù)HD高清標準和4K的來(lái)臨,在CPU上運行應用程序嚴重受制于其軟件的熱封裝屬性——如果只是增加多個(gè)內核,以高頻運行,會(huì )造成晶體管過(guò)熱并關(guān)機則并不可行。這個(gè)問(wèn)題亟待解決。 PowerVRSGX GPU以70 fps對Full HD(1080p)實(shí)時(shí)視頻運行飽和過(guò)濾器 還好可以利用這個(gè)解決方案:所有先進(jìn)應用處理器包含多個(gè)計算引擎,如GPU、DSP和ISP(圖像合成處理器),可有效地執行所需的成像任務(wù)。因此,開(kāi)發(fā)人員正在研究異構計算,也就是將這些塊集成,在有限的電源和散熱預算范圍內提供高性能。 全球最具代表性的手機上都會(huì )找到PowerVR GPU,以低功率提供世界一流的3D圖形。使用非常靈活的架構,集成業(yè)界領(lǐng)先的PowerVR GPU可部署高效計算引擎來(lái)加速圖像和視頻處理任務(wù)。應用程序會(huì )通過(guò)OpenCL EP(嵌入式文件)API要求GPU的成像內核提供視覺(jué)計算的最佳精度,同時(shí)維持最低功耗。 這就提出一個(gè)問(wèn)題,即如何共享各種計算引擎之間的圖像數據,無(wú)需冗余運算。這篇博客的其余部分將說(shuō)明如何能夠利用最近向開(kāi)發(fā)人員發(fā)布Imagination開(kāi)發(fā)的工具。 共享所面臨的挑戰——確保零拷貝 App如果要在設計上充分發(fā)揮異構計算的優(yōu)勢就要求用于編程不同計算引擎的API之間有高效的互操作性。例如在實(shí)時(shí)電話(huà)會(huì )議應用程序當中,可能首先由GPU檢查相機捕獲的輸入幀決定臉部位置,以及各項特征(如眼睛、嘴唇、鼻子和其它部位),將這些坐標傳送到CPU進(jìn)行分析。然后,CPU可以判斷可由GPU使用的一組圖像過(guò)濾器,如除去瑕疵或皺紋,GPU生成變換幀,輸入至視頻編解碼器。在這種情況下,四個(gè)不同的系統組件分別需要訪(fǎng)問(wèn)內存中同樣的圖像數據。 到現在為止,市場(chǎng)上所有OpenCL實(shí)現創(chuàng )建圖像數據的后臺復本,在相機和GPU之間傳遞。這個(gè)操作徒勞地增加了系統內存流量,降低性能和功耗,影響(在某些情況下消除)卸載GPU任務(wù)的效能。過(guò)去一年來(lái)Imagination一直與合作伙伴合作致力于消除這一障礙,使用OpenCL提高基于相機應用程序效能。 如何解決這個(gè)難題? 我們開(kāi)發(fā)了一套擴展,使共享同一系統內存的多個(gè)組件可共享圖像。這些擴展基于Khronos EGL圖像,提供多個(gè)Khronos API之間的接口、本地平臺窗口系統,并處理與綁定和同步相關(guān)的問(wèn)題。 許多相機生成YUV格式的圖像數據,其中的Y和UV數據被存儲在分開(kāi)的平面。PowerVR SGXGPU在其固定功能紋理化硬件上可以自動(dòng)執行按向量運算執行色彩空間轉換將YUV(NV21)轉換為RGB,并處理各像素,對R、G和B像素值進(jìn)行有效的并行計算。OpenCL內核的結構與以下示例類(lèi)似。輸入和輸出圖像被表示為OpenCL的圖像數據的類(lèi)型,各個(gè)像素被處理為float4數據類(lèi)型,與本地SGX向量寬度匹配。 /* Generic 3x3 linear filter kernel. */ __kernel void convolve3x3(__read_only image2d_tsrcImage, __write_only image2d_t dstImage, sampler_t sampler, __constant float *kVals, float normalizationVal) { int2 coords = (int2)(get_global_id(0),get_global_id(1)); float4 colour; float4 colours[9]; colours[0] = read_imagef(srcImage, sampler, coords +(int2)(-1, -1)); colours[1] = read_imagef(srcImage, sampler, coords +(int2)( 0, -1)); colours[2] = read_imagef(srcImage, sampler, coords +(int2)( 1, -1)); colours[3] = read_imagef(srcImage, sampler, coords +(int2)(-1, 0)); colours[4] = read_imagef(srcImage, sampler, coords +(int2)( 0, 0)); colours[5] = read_imagef(srcImage, sampler, coords +(int2)( 1, 0)); colours[6] = read_imagef(srcImage, sampler, coords +(int2)(-1, 1)); colours[7] = read_imagef(srcImage, sampler, coords +(int2)( 0, 1)); colours[8] = read_imagef(srcImage, sampler, coords +(int2)( 1, 1)); colour = colours[0] * kVals[0] + colours[1] * kVals[1] +colours[2] * kVals[2]; colour += colours[3] * kVals[3] + colours[4] * kVals[4] +colours[5] * kVals[5]; colour += colours[6] * kVals[6] + colours[7] * kVals[7] +colours[8] * kVals[8]; colour /= normalizationVal; write_imagef(dstImage, coords, colour); } 完整的系統解決方案如下圖所示。使用Khronos擴展OES_EGL_image_external將相機YUV數據轉換成RGB,然后利用PowerVR的零拷貝功能,直接訪(fǎng)問(wèn)OpenCL的相機紋理數據,不需要從OpenGLES到OpenCL執行中間復制。濾波后的輸出圖像插入另一個(gè)零拷貝紋理,可以用OpenGLES寫(xiě)入2D或3D畫(huà)面。 多個(gè)系統組件之間的零拷貝數據 執行OpenCL視頻處理 Imagination在最近一系列視頻中演示以Full HD(1080p)運行的視頻率計算攝影。本演示演示如何為Samsung Exynos5410,編寫(xiě)“類(lèi)似Instagram”應用程序,在三星Galaxy S4 i9500和魅族MX3智能手機以及HardkernelODROID-XU開(kāi)發(fā)板上均能找到這類(lèi)應用程序,與在CPU(4-7幀)上運行圖像處理任務(wù)相比,可利用PowerVRGPU實(shí)現最佳性能,同時(shí)降低功耗。 為了使開(kāi)發(fā)人員能夠創(chuàng )建自有的高性能GPU計算應用程序,Imagination即將公布GPU計算開(kāi)發(fā)計劃,注冊用戶(hù)將有機會(huì )訪(fǎng)問(wèn)PowerVRGPU計算SDK和PowerVR GPU編程指南。在這篇博客中所討論的擴展完全受到OpenCL的驅動(dòng)程序的支持,可接入低開(kāi)銷(xiāo)Hardkernel ORDOID -XU板。此外,如果希望充分利用Imagination PowerVR6系列GPU計算功能可使用PVRTune工具,其中將包括OpenCL和API的計時(shí)數據的增強配置文件,使開(kāi)發(fā)人員能夠編寫(xiě)出更好的應用程序。 |