使用PowerVR SGX GPU實(shí)現(xiàn)類似Instagram的計(jì)算攝影

　　自從首部拍照手機(jī)問世后，手機(jī)就成為視覺導(dǎo)向設(shè)備，今天用戶產(chǎn)生的內(nèi)容已成為從網(wǎng)絡(luò)利用率到創(chuàng)建應(yīng)用程序的主要推手，然而，開發(fā)人員們最近開始對(duì)圖像數(shù)據(jù)進(jìn)行高級(jí)處理，以創(chuàng)建一系列新的用戶體驗(yàn)，這些體驗(yàn)從智能視覺系統(tǒng)、增強(qiáng)現(xiàn)實(shí)應(yīng)用程序到虛擬制作等等。不過這些創(chuàng)新技術(shù)的爆發(fā)受到手機(jī)數(shù)據(jù)處理量的限——這當(dāng)中存在很大挑戰(zhàn)。

　　計(jì)算攝影：移動(dòng)計(jì)算的主要應(yīng)用之一

　　成像一直以來是CPU的主任務(wù)，對(duì)于靜止圖像和小格式視頻來說CPU處理得很好，但隨著HD高清標(biāo)準(zhǔn)和4K的來臨，在CPU上運(yùn)行應(yīng)用程序嚴(yán)重受制于其軟件的熱封裝屬性——如果只是增加多個(gè)內(nèi)核，以高頻運(yùn)行，會(huì)造成晶體管過熱并關(guān)機(jī)則并不可行。這個(gè)問題亟待解決。

　　PowerVRSGX GPU以70 fps對(duì)Full HD(1080p)實(shí)時(shí)視頻運(yùn)行飽和過濾器

　　還好可以利用這個(gè)解決方案：所有先進(jìn)應(yīng)用處理器包含多個(gè)計(jì)算引擎，如GPU、DSP和ISP(圖像合成處理器)，可有效地執(zhí)行所需的成像任務(wù)。因此，開發(fā)人員正在研究異構(gòu)計(jì)算，也就是將這些塊集成，在有限的電源和散熱預(yù)算范圍內(nèi)提供高性能。

　　全球最具代表性的手機(jī)上都會(huì)找到PowerVR GPU，以低功率提供世界一流的3D圖形。使用非常靈活的架構(gòu)，集成業(yè)界領(lǐng)先的PowerVR GPU可部署高效計(jì)算引擎來加速圖像和視頻處理任務(wù)。應(yīng)用程序會(huì)通過OpenCL EP(嵌入式文件)API要求GPU的成像內(nèi)核提供視覺計(jì)算的最佳精度，同時(shí)維持最低功耗。

　　這就提出一個(gè)問題，即如何共享各種計(jì)算引擎之間的圖像數(shù)據(jù)，無需冗余運(yùn)算。這篇博客的其余部分將說明如何能夠利用最近向開發(fā)人員發(fā)布Imagination開發(fā)的工具。

　　共享所面臨的挑戰(zhàn)——確保零拷貝

　　App如果要在設(shè)計(jì)上充分發(fā)揮異構(gòu)計(jì)算的優(yōu)勢就要求用于編程不同計(jì)算引擎的API之間有高效的互操作性。例如在實(shí)時(shí)電話會(huì)議應(yīng)用程序當(dāng)中，可能首先由GPU檢查相機(jī)捕獲的輸入幀決定臉部位置，以及各項(xiàng)特征(如眼睛、嘴唇、鼻子和其它部位)，將這些坐標(biāo)傳送到CPU進(jìn)行分析。然后，CPU可以判斷可由GPU使用的一組圖像過濾器，如除去瑕疵或皺紋，GPU生成變換幀，輸入至視頻編解碼器。在這種情況下，四個(gè)不同的系統(tǒng)組件分別需要訪問內(nèi)存中同樣的圖像數(shù)據(jù)。

　　到現(xiàn)在為止，市場上所有OpenCL實(shí)現(xiàn)創(chuàng)建圖像數(shù)據(jù)的后臺(tái)復(fù)本，在相機(jī)和GPU之間傳遞。這個(gè)操作徒勞地增加了系統(tǒng)內(nèi)存流量，降低性能和功耗，影響(在某些情況下消除)卸載GPU任務(wù)的效能。過去一年來Imagination一直與合作伙伴合作致力于消除這一障礙，使用OpenCL提高基于相機(jī)應(yīng)用程序效能。

　　如何解決這個(gè)難題?

　　我們開發(fā)了一套擴(kuò)展，使共享同一系統(tǒng)內(nèi)存的多個(gè)組件可共享圖像。這些擴(kuò)展基于Khronos EGL圖像，提供多個(gè)Khronos API之間的接口、本地平臺(tái)窗口系統(tǒng)，并處理與綁定和同步相關(guān)的問題。

　　許多相機(jī)生成YUV格式的圖像數(shù)據(jù)，其中的Y和UV數(shù)據(jù)被存儲(chǔ)在分開的平面。PowerVR SGXGPU在其固定功能紋理化硬件上可以自動(dòng)執(zhí)行按向量運(yùn)算執(zhí)行色彩空間轉(zhuǎn)換將YUV(NV21)轉(zhuǎn)換為RGB，并處理各像素，對(duì)R、G和B像素值進(jìn)行有效的并行計(jì)算。OpenCL內(nèi)核的結(jié)構(gòu)與以下示例類似。輸入和輸出圖像被表示為OpenCL的圖像數(shù)據(jù)的類型，各個(gè)像素被處理為float4數(shù)據(jù)類型，與本地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);

　　}