自從首部拍照手機問世后，手機就成為視覺導向設備，今天用戶產生的內容已成為從網絡利用率到創建應用程序的主要推手，然而，開發人員們最近開始對圖像數據進行高級處理，以創建一系列新的用戶體驗，這些體驗從智能視覺系統、增強現實應用程序到虛擬制作等等。不過這些創新技術的爆發受到手機數據處理量的限——這當中存在很大挑戰。

　　計算攝影：移動計算的主要應用之一

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

　　PowerVRSGX GPU以70 fps對Full HD(1080p)實時視頻運行飽和過濾器

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

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

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

　　共享所面臨的挑戰——確保零拷貝

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

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

　　如何解決這個難題?

　　我們開發了一套擴展，使共享同一系統內存的多個組件可共享圖像。這些擴展基于Khronos EGL圖像，提供多個Khronos API之間的接口、本地平臺窗口系統，并處理與綁定和同步相關的問題。

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

　　}