使用PowerVR SGX GPU實(shí)現類(lèi)似Instagram的計算攝影

　　自從首部拍照手機問(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);

　　}