使用PowerVR SGX GPU实现类似Instagram的计算摄影

　　自从首部拍照手机问世后，手机就成为视觉导向设备，今天用户产生的内容已成为从网络利用率到创建应用程序的主要推手，然而，开发人员们最近开始对图像数据进行高级处理，以创建一系列新的用户体验，这些体验从智能视觉系统、增强现实应用程序到虚拟制作等等。不过这些创新技术的爆发受到手机数据处理量的限——这当中存在很大挑战。

　　计算摄影：移动计算的主要应用之一

　　成像一直以来是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);

　　}