65.9K
CodeProject 正在变化。 阅读更多。
Home

将图像处理卸载到 Android 应用程序的编程

2015年10月1日

CPOL

21分钟阅读

viewsIcon

20287

本文通过一个示例安卓应用,演示了如何使用 OpenCL™ 和 RenderScript 编程语言卸载图像处理任务。

英特尔® 开发人员专区提供用于跨平台应用开发的工具和操作指南、平台和技术信息、代码示例以及同行专业知识,旨在帮助开发人员进行创新并取得成功。加入我们的安卓物联网英特尔® 实感™ 技术Windows社区,下载工具、访问开发套件、与志同道合的开发人员分享想法,并参与黑客马拉松、竞赛、路演和本地活动。

1. 引言

本文将通过一个示例安卓应用,介绍如何使用 OpenCL™ 和 RenderScript 编程语言来卸载图像处理任务。这些编程语言旨在利用高度并行的图形硬件(着色器)来处理大数据集和高度重复的任务。虽然您可以使用其他语言在安卓应用中卸载图像处理,但本文将展示 OpenCL 和 RenderScript 示例代码,用于开发应用基础架构和图像处理算法代码。文中还将展示 OpenCL API 包装类,该类用于促进图像处理算法的 OpenCL 编程和执行。OpenCL API 包装类的源代码可免费供任何人使用。

您应熟悉 OpenCL、RenderScript 和安卓编程概念,因为本文仅涵盖卸载图像处理或媒体生成计算的说明。您还应拥有一台已配备、启用和配置为运行 OpenCL 的安卓设备(有关安卓设备安装,请参阅英特尔® OpenCL SDK)。

注意:虽然还有其他语言和技术可用于卸载图像处理或媒体生成,但本文的目标是仅突出代码差异。未来计划发布的文章将突出 OpenCL 和 RenderScript 在 GPU 上执行时的性能差异。

1.1 应用程序 UI 设计

在示例应用程序中,UI 有三个单选按钮,用户可以在 RenderScript、OpenCL 或原生代码之间快速切换应用程序执行。另一个菜单设置允许用户选择是在 CPU 还是 GPU 上运行 OpenCL。菜单还为用户提供了可运行的已实现效果列表,因此他们可以选择要运行的效果。选择设备仅适用于 OpenCL(不适用于 RenderScript 或原生代码)。英特尔® x86 平台在 CPU 和 GPU 上都包含 OpenCL 运行时支持。

以下是主 UI 的屏幕截图,显示了 OpenCL 处理等离子效果的一个版本。示例应用程序 UI 显示了运行 OpenCL、RenderScript 或原生代码时的性能结果。

性能指标包括每秒帧数 (fps)、帧渲染和效果计算的耗时。性能指标在下面的屏幕截图中突出显示。

请注意,屏幕截图上显示的性能数字是示例指标;实际性能指标将因设备而异。

1.1 API 和 SDK

除了 ADT(Android 开发工具,其中还包括 Android SDK),用于编写示例应用程序的主要基于 Android 的 API 是 RenderScript 和英特尔® SDK for OpenCL for Android 应用程序。

英特尔 OpenCL™ SDK 基于并遵循 OpenCL™ 规范,该规范是一个开放、免版税的跨平台编程标准。有关更多详细信息,请参阅 Khronos 网站上的 OpenCL™ 标准。

RenderScript 最初在 2.2 ADT(API 级别 8)中提供,是一个用于在 Android 上运行计算密集型任务的框架。RenderScript 主要用于数据并行计算,但串行计算工作负载也能从中受益。有关更多信息,请参阅 Android 开发者网站

Google 开源存储库中提供的最新 ADT 包含需要导入的适当软件包,以使用 RenderScript、JNI (Java* Native Interface) 和运行时 API。有关 OpenCL 的设置、配置和运行时,请参阅这篇适用于 Android OS 的 OpenCL 开发文章。有关其他编程细节,请参阅 RenderScriptOpenCL

1.3 基础设施代码

基础设施代码由“主”活动和辅助函数组成。本节重点介绍用于设置 UI、选择要运行的效果和语言技术,以及对于 OpenCL,选择要使用的计算设备的辅助函数和代码。

虽然实现了几个辅助函数来集成用户选择命令,但这里只突出显示了两个

辅助函数 backgroundThread() 启动一个线程,该线程定期调用 step process 函数来处理图像效果。此函数中使用的代码和功能取自 RenderScript 入门文章中发布的另一个示例应用程序,您可以在此处(PDF)找到更多详细信息。

processStep() 函数由 backgroundThread() 调用,用于处理和运行图像效果。该函数依赖于一个单选按钮回调函数来确定要使用的语言。processStep() 函数调用适当的方法来使用 OpenCL、RenderScript 或纯原生 C/C++ 代码处理图像效果。由于此代码在后台线程上运行,用户只需单击或触摸单选按钮即可选择要运行的语言,即使在处理效果时也是如此。应用程序会动态切换以执行给定图像效果的适当步骤渲染函数。

// The processStep() method runs in a separate (background) thread.
private void processStep() {
	try {
		switch (this.type.getCheckedRadioButtonId()) {
		case R.id.type_renderN:
			oclFlag = 0; // OpenCL is OFF
			stepRenderNative();
			break;
		case R.id.type_renderOCL:
			oclFlag = 1; // OpenCL is ON
			stepRenderOpenCL();
			break;
		case R.id.type_renderRS:
		      oclFlag = 0; // OpenCL is OFF
			stepRenderScript();
			break;
		default:
			return;
		}
	} catch (RuntimeException ex) {
		// Handle exception as appropriate and log error
		Log.wtf("Android Image Processing", "render failed", ex);
	}
}

1.4 原生函数的 Java 定义

示例应用程序实现了一个 NativeLib 类,该类主要定义通过 JNI 调用原生功能以处理给定效果的函数。例如,示例应用程序实现了三种效果:等离子、棕褐色和单色。因此,该类定义了 renderPlasma(…)renderSepia(…)renderMonoChrome(…) 函数。这些 Java 函数作为通过 JNI 运行原生或 OpenCL 功能的入口点。

JNI 函数执行 C/C++ 代码或设置并执行实现图像效果的 OpenCL 程序。该类使用 Android 位图和 AssetManager 包。BitMap 对象用于传递和返回正在处理的图像或媒体的数据。应用程序依赖于 AssetManager 对象来访问 OpenCL 文件(即 sepia.cl),其中定义了 OpenCL 内核。

以下是实际的 NativeLib Java 类定义。包含 //TODO 注释是为了说明应用程序可以轻松扩展以实现额外的图像效果。

package com.example.imageprocessingoffload;
import android.content.res.AssetManager;
import android.graphics.Bitmap;

public class NativeLib
{
    // Implemented in libimageeffects.so
    public static native void renderPlasma(Bitmap bitmapIn, int renderocl, long time_ms, String eName, int devtype, AssetManager mgr);
       
    public static native void renderMonoChrome(Bitmap bitmapIn, Bitmap bitmapOut, int renderocl, long time_ms, String eName, int simXtouch, int simYtouch, int radHi, int radLo, int devtype, AssetManager mgr);

    public static native void renderSepia(Bitmap bitmapIn, Bitmap bitmapOut, int renderocl, long time_ms, String eName, int simXtouch, int simYtouch, int radHi, int radLo, int devtype, AssetManager mgr);
 
    //TODO public static native <return type> render<Effectname>(…);
    
    //load actual native library
    static {
        System.loadLibrary("imageeffectsoffloading");
    }
}

请注意,Android AssetManagerBitMap 对象作为图像输入和图像结果传递给原生代码。原生代码使用 AssetManager 对象来访问定义 OpenCL 内核的 CL 文件。BitMap 对象用于使像素数据可供原生代码计算并生成图像结果。

UI 参数 deviceType 用于指示是否在 CPU 或 GPU 上执行 OpenCL。Android 系统必须配置并能够在这两种设备上运行 OpenCL。现代英特尔® Atom™ 和英特尔® 酷睿™ 处理器可以在 CPU 和集成显卡处理器或 GPU 上运行 OpenCL。

参数 eName 用于指示要编译和运行哪个 OpenCL 内核。尽管示例应用程序为每个图像效果实现了一个 JNI 函数,但这似乎没有必要。然而,可以在单个 CL 文件和/或 JNI 函数中定义多个相关的图像效果。在这种情况下,eName 将用于编译和加载适当的 CL 程序和/或内核。

参数 renderocl 用作一个标志,指示是运行 OpenCL 还是原生 C/C++ 代码。此标志仅在用户选择 OpenCL 单选按钮时设置;否则,它保持未设置状态。

参数 time_ms 用于传递时间戳(毫秒),用于计算性能指标。在等离子效果中,时间戳用于计算等离子效果步进。

其他参数特定于图像效果算法,用于从图像中心径向渲染效果。例如,simXtouchsimYtouchradLoradHi 参数以及宽度和高度用于计算和显示单色和棕褐色效果的径向进度。

1.5 运行原生代码(C 或 OpenCL)的定义和资源

本节包括示例应用程序中实现的每种效果的 JNI 本机函数定义。如前所述,每种效果使用一个函数,以简化解释并说明用于通过 OpenCL 卸载图像效果处理的功能元素。C 或串行代码被引用,并且还包含代码片段,希望未来版本的示例应用程序将用于评估这些语言技术之间的性能。

JNI 函数与 Java 原生函数具有 1:1 的关系。因此,正确声明和定义 JNI 对应函数非常重要。Java SDK 包含 javah 工具,该工具可帮助生成正确且精确的 JNI 函数声明。强烈建议使用此工具,以避免代码编译正确但在运行时产生错误可能导致的问题。

以下是示例应用程序中“图像效果卸载”的 JNI 函数。JNI 函数签名由 javah 工具实用程序生成。

// Defines new JNI entry function signatures
#ifndef _Included_com_example_imageprocessingoffload_NativeLib
#define _Included_com_example_imageprocessingoffload_NativeLib
#ifdef __cplusplus
extern "C" {
#endif
/*
 * Class:     com_example_imageprocessingoffload_NativeLib
 * Method:    renderPlasma
 * Signature: (Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String;
 */
JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderPlasma
  (JNIEnv *, jclass, jobject, jint, jlong, jstring, jint, jobject);

/*
 * Class:     com_example_imageprocessingoffload_NativeLib
 * Method:    renderMonoChrome
 * Signature: (Landroid/graphics/Bitmap;Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String;
 */
JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderMonoChrome
  (JNIEnv *, jclass, jobject, jobject, jint, jlong, jstring, jint, jint, jint, jint, jint, jobject);

/*
 * Class:     com_example_imageprocessingoffload_NativeLib
 * Method:    renderSepia
 * Signature: (Landroid/graphics/Bitmap;Landroid/graphics/Bitmap;IJLjava/lang/String;)Ljava/lang/String;
 */
JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderSepia
  (JNIEnv *, jclass, jobject, jobject, jint, jlong, jstring, jint, jint, jint, jint, jint, jobject);
}
#endif

javah 工具可以生成正确的 JNI 函数签名;但是,定义 Java 原生函数的类必须已在您的 Android 应用程序项目中编译。如果需要生成头文件,javah 命令可以如下使用

{javahLocation} -o {outputFile} -classpath {classpath} {importName}

对于示例应用程序,函数签名生成如下

javah -o junk.h -classpath bin\classes com.example.imageprocessingoffloading.NativeLib

然后,junk.h 中的 JNI 函数签名被添加到 *imageeffects.cpp* 中,该文件具有设置和运行 OpenCL 或 C 代码的功能。接下来,我们分配资源以能够运行 OpenCL 或原生代码以实现等离子、单色和棕褐色效果。

1.5.1 等离子效果

Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(…) 函数是执行等离子效果的 OpenCL 或原生代码的入口代码。函数 startPlasmaOpenCL(…)runPlasmaOpenCL(…)runPlasmaNative(…) 位于 *imageeffects.cpp* 代码之外,并定义在单独的 *plasmaEffect.cpp* 源文件中。作为参考,您可以在 OpenCL 包装类代码下载中找到 *plasmaEffect.cpp* 源文件。

renderPlasma(…) 入口函数利用 OpenCL 包装类查询 Android 设备系统对 OpenCL 的支持。它调用包装类函数 ::initOpenCL(…) 初始化 OpenCL 环境。设备类型将 CPU 或 GPU 作为设备传递以创建 OpenCL 上下文。Android 资产管理器使用 ceName 参数识别并加载用于编译内核代码的 CL 文件。

如果 OpenCL 环境成功设置,renderPlasma(…) 入口函数将调用 startPlasmaOpenCL() 函数来分配 OpenCL 资源并开始执行等离子 OpenCL 内核。请注意,gOCL 是一个全局变量,它保存 OpenCL 包装类的对象实例。gOCL 变量对所有 JNI 入口函数都可见。这样,OpenCL 环境可以由任何已编程的效果初始化。

等离子效果不使用图像,屏幕上渲染的媒体由编程算法生成。bitmapIn 参数是一个 BitMap 对象,它保存由等离子效果生成的媒体。在 startPlasma(…) 函数中传递的 pixels 参数被映射到位图纹理,并由原生或 OpenCL 内核代码用于读取和写入用于在屏幕上渲染纹理的像素数据。再次强调,assetManager 对象用于访问包含等离子效果 OpenCL 内核的 CL 文件。

JNIEXPORT void Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(JNIEnv * env, jclass, jobject bitmapIn, jint renderocl, jlong time_ms, jstring ename, jint devtype, jobject assetManager) {

JNIEXPORT void Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(JNIEnv * env, jclass, jobject bitmapIn, jint renderocl, jlong time_ms, jstring ename, jint devtype, jobject assetManager) { 
… // code omitted to simplify	
    
    // code locks mem for BitMapIn and sets “pixels” pointer that is passed to OpenCL or Native functions.  
    ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixels); 
 
… // code omitted to simplify

  If OCL not initialized
     AAssetManager *amgr = AAssetManager_fromJava(env, assetManager);
     gOCL.initOpenCL(clDeviceType, ceName, amgr);
     startPlasmaOpenCL((cl_ushort *) pixels, infoIn.height, infoIn.width, (float) time_ms, ceName, cpinit);
 else
     runPlasmaOpenCL(infoIn.width, infoIn.height, (float) time_ms, (cl_ushort *) pixels);
… // code omitted
}

startPlasmaOpenCL(…) 外部函数生成并填充包含等离子效果所需数据的 PaletteAngles 缓冲区。为了开始运行等离子 OpenCL 内核,该函数依赖于 OpenCL 命令队列、上下文和内核,这些都在包装类的成员数据中定义。

runPlasmaOpenCL(…) 函数持续运行等离子 OpenCL 内核。一旦 OpenCL 内核启动,就会使用一个单独的函数,后续的内核执行只需要一个新的时间戳值作为输入。下一次内核运行迭代只需要发送时间戳值的内核参数,因此需要一个单独的函数。

extern int startPlasmaOpenCL(cl_ushort* pixels, cl_int height, cl_int width, cl_float ts, const char* eName, int inittbl);
extern int runPlasmaOpenCL(int width, int height, cl_float ts, cl_ushort *pixels);
extern void runPlasmaNative( AndroidBitmapInfo*  info, void*  pixels, double  t, int inittbl );

runPlasmaNative(…) 函数包含用 C 代码编写的等离子算法逻辑。inittbl 参数用作布尔值,表示是否需要生成等离子效果所需的 PaletteAngles 数据。等离子效果的 OpenCL 内核代码可以在 *plasmaEffect.cpp* 源文件中找到。

#define FBITS		16
#define FONE		(1 << FBITS)
#define FFRAC(x)	((x) & ((1 << FBITS)-1))
#define FIXED_FROM_FLOAT(x)  ((int)((x)*FONE))

/* Color palette used for rendering plasma */
#define  PBITS   8
#define  ABITS   9
#define  PSIZE   (1 << PBITS)
#define  ANGLE_2PI (1 << ABITS)
#define  ANGLE_MSK (ANGLE_2PI - 1)
 
#define  YT1_INCR  FIXED_FROM_FLOAT(1/100.0f)
#define  YT2_INCR  FIXED_FROM_FLOAT(1/163.0f)
#define  XT1_INCR  FIXED_FROM_FLOAT(1/173.0f)
#define  XT2_INCR  FIXED_FROM_FLOAT(1/242.0f)
 
#define  ANGLE_FROM_FIXED(x)	((x) >> (FBITS - ABITS)) & ANGLE_MSK

ushort pfrom_fixed(int x, __global ushort *palette)
{
    if (x < 0) x = -x;
    if (x >= FONE) x = FONE-1;
    int  idx = FFRAC(x) >> (FBITS - PBITS);
    return palette[idx & (PSIZE-1)];
}
 
__kernel
void plasma(__global ushort *pixels, int height, int width, float t, __global ushort *palette, __global int *angleLut)
{
    int yt1 = FIXED_FROM_FLOAT(t/1230.0f); 
    int yt2 = yt1;
    int xt10 = FIXED_FROM_FLOAT(t/3000.0f);
    int xt20 = xt10;
 
    int x = get_global_id(0);
    int y = get_global_id(1);
    int tid = x+y*width;
 
    yt1 += y*YT1_INCR;
    yt2 += y*YT2_INCR;
 
    int base = angleLut[ANGLE_FROM_FIXED(yt1)] + angleLut[ANGLE_FROM_FIXED(yt2)];
    int xt1 = xt10;
    int xt2 = xt20;
 
    xt1 += x*XT1_INCR;
    xt2 += x*XT2_INCR;
 
    int ii = base + angleLut[ANGLE_FROM_FIXED(xt1)] + angleLut[ANGLE_FROM_FIXED(xt2)];
    pixels[tid] = pfrom_fixed(ii/4, palette);
}
The RenderScript kernel code for the plasma effect:

#pragma version(1)
#pragma rs java_package_name(com.example.imageprocessingoffload)

rs_allocation *gPalette;
rs_allocation *gAngles;
rs_script gScript;
float ts;
int gx;
int gy;

static int32_t intFromFloat(float xfl) {
      return (int32_t)((xfl)*(1 << 16));
}
const float YT1_INCR = (1/100.0f);
const float YT2_INCR = (1/163.0f);
const float XT1_INCR = (1/173.0f);
const float XT2_INCR = (1/242.0f);

static uint16_t pfrom_fixed(int32_t dx) {
    unsigned short *palette = (unsigned short *)gPalette;
    uint16_t ret;
    if (dx < 0)  dx = -dx;
    if (dx >= (1 << 16))  dx = (1 << 16)-1;   

    int  idx = ((dx & ((1 << 16)-1)) >> 8);
    ret = palette[idx & ((1<<8)-1)];
    return ret;
}

uint16_t __attribute__((kernel)) root(uint16_t in, uint32_t x, uint32_t y) {
    unsigned int *angles = (unsigned int *)gAngles;
    uint32_t out = in;
    int yt1 = intFromFloat(ts/1230.0f); 

    int yt2 = yt1;
    int xt10 = intFromFloat(ts/3000.0f);
    int xt20 = xt10;
    
    int y1 = y*intFromFloat(YT1_INCR);
    int y2 = y*intFromFloat(YT2_INCR);
    yt1 = yt1 + y1;
    yt2 = yt2 + y2;
    
    int a1 = (yt1 >> 7) & ((1<<9)-1);
    int a2 = (yt2 >> 7) & ((1<<9)-1);
    int base = angles[a1] + angles[a2];
    
    int xt1 = xt10;
    int xt2 = xt20;
    xt1 += x*intFromFloat(XT1_INCR);
    xt2 += x*intFromFloat(XT2_INCR);
	
    a1 = (xt1 >> (16-9)) & ((1<<9)-1);
    a2 = (xt2 >> (16-9)) & ((1<<9)-1);
    int ii = base + angles[a1] + angles[a2];
	
   out = pfrom_fixed(ii/4);
   return out;
}
void filter(rs_script gScript, rs_allocation alloc_in, rs_allocation alloc_out) {
    //rsDebug("Inputs TS, X, Y:", ts, gx, gy);
    rsForEach(gScript, alloc_in, alloc_out);
}

1.5.2 单色效果

Java_com_example_imageprocessingoffload_NativeLib_renderMonochrome(…) 函数是执行单色处理的 OpenCL 或原生代码的入口代码。函数 executeMonochromeOpenCL(…)executeMonochromeNative(…) 位于 *imageeffects.cpp* 代码之外,并定义在单独的源文件中。与等离子效果一样,此入口函数也利用 OpenCL 包装类查询 Android 设备系统对 OpenCL 的支持,并调用函数 ::initOpenCL(…) 初始化 OpenCL 环境。

以下两行代码简单地 extern(或使 NDK 编译器可见)executeMonochromeOpenCL(…)executeMonochromeNative(…) 函数的函数签名。这些行是必需的,因为这些函数定义在单独的源文件中。

extern int executeMonochromeOpenCL(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight);
extern int executeMonochromeNative(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight);

与等离子效果不同,此效果使用输入图像和输出图像。bitmapInbitmapOut 都被分配为 ARGB_8888 位图,并且都映射到 cl_uchar4 向量的 CL 缓冲区。请注意,pixelsInpixelsOut 被强制转换,因为 OpenCL 必须将 BitMap 对象映射到 cl_uchar4 向量的缓冲区。

JNIEXPORT void JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderMonochrome(JNIEnv * env, jclass obj, jobject bitmapIn, jobject bitmapOut, jint renderocl, jlong time_ms, jstring ename, jint xto, jint yto, jint radHi, jint radLo, jint devtype, jobject assetManager)  {

  … // code omitted for simplification

   // code locks mem for BitMapIn and sets “pixelsIn” pointer that is passed to OpenCL or Native functions.  
   ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixelsIn); 

   // code locks mem for BitMapOut and sets “pixelsOut” pointer that is passed to OpenCL or Native functions.  
   ret = AndroidBitmap_lockPixels(env, bitmapOut, &pixelsOut); 

 … // code omitted for simplification
 If OpenCL
   If OCL not initialized
     AAssetManager *amgr = AAssetManager_fromJava(env, assetManager);
     gOCL.initOpenCL(clDeviceType, ceName, amgr);
   else
     executeMonochromeOpenCL((cl_uchar4*) pixelsIn,(cl_uchar4*) pixelsOut, radiHi, radiLo, xt, yt, infoIn.width, infoIn.height);
    // end of OCL initialized
else
   executeMochromeNative((cl_uchar4*) pixelsIn,(cl_uchar4*) pixelsOut, radiHi, radiLo, xt, yt, infoIn.width, infoIn.height);
// End of OpenCL
… // code omitted
}

当调用 executeMonochromeOpenCL(…) 时,该函数将 pixelsInpixelsOut 强制转换为 cl_uchar4 缓冲区并传递。该函数使用 OpenCL API 创建缓冲区和其他适当的资源。它设置内核参数并排队必要的命令来执行 OpenCL 内核。由 pixelsIn 指向的图像输入缓冲区被分配为只读缓冲区。内核代码使用 pixelsIn 指针获取传入的像素数据。像素数据被内核算法用于将传入图像转换为单色图像。输出缓冲区是读写缓冲区,它保存图像结果并由 pixelsOut 指向。有关 OpenCL 的更多详细信息,请参阅英特尔的编程和优化指南

executeMonochromeNative(…) 函数包含用 C 代码编写的单色算法。该算法是基本的,由一个外循环(y 循环)和一个内循环(x 循环)组成,用于计算像素数据,其结果存储在由 pixelsOut 指向的 dstImage 中。由 pixlesIn 指向的 srcImage 用于解引用输入像素数据,以便算法公式将其转换为单色像素。

单色效果的 OpenCL 内核代码

constant uchar4 cWhite = {1.0f, 1.0f, 1.0f, 1.0f};
constant float3 channelWeights = {0.299f, 0.587f, 0.114f};
constant float saturationValue = 0.0f;

__kernel void mono (__global uchar4 *in, __global uchar4 *out, int4 intArgs, int width) {
    int x = get_global_id(0);
    int y = get_global_id(1);
   
    int xToApply = intArgs.x;
    int yToApply = intArgs.y;
    int radiusHi = intArgs.z;
    int radiusLo = intArgs.w;
    int tid = x + y * width;
    uchar4 c4 = in[tid];
    float4 f4 = convert_float4 (c4);
    int xRel = x - xToApply;
    int yRel = y - yToApply;
    int polar = xRel*xRel + yRel*yRel;
   
    if (polar > radiusHi || polar < radiusLo)   {
        if (polar < radiusLo)   {
            float4 outPixel = dot (f4.xyz, channelWeights);
            outPixel = mix ( outPixel, f4, saturationValue);
            outPixel.w = f4.w;
            out[tid] = convert_uchar4_sat_rte (outPixel); 
        }
        else  {
            out[tid] = convert_uchar4_sat_rte (f4);
        }
    }
    else   {
         out[tid] = convert_uchar4_sat_rte (cWhite);
    }
}

The RenderScript kernel code for the monochrome effect:

#pragma version(1)
#pragma rs java_package_name(com.example.imageprocessingoffload)

int radiusHi;
int radiusLo;
int xToApply;
int yToApply;

const float4 gWhite = {1.f, 1.f, 1.f, 1.f};
const float3 channelWeights = {0.299f, 0.587f, 0.114f};
float saturationValue = 0.0f;

uchar4 __attribute__((kernel)) root(const uchar4 in, uint32_t x, uint32_t y)
{
    float4 f4 = rsUnpackColor8888(in);
    int xRel = x - xToApply;
    int yRel = y - yToApply;
    int polar = xRel*xRel + yRel*yRel;
    uchar4 out;
    
    if(polar > radiusHi || polar < radiusLo) {
        if(polar < radiusLo) {
            float3 outPixel = dot(f4.rgb, channelWeights);
            outPixel = mix( outPixel, f4.rgb, saturationValue);
            out = rsPackColorTo8888(outPixel);
        }
        else {
            out = rsPackColorTo8888(f4);
        }
    }
    else {
         out = rsPackColorTo8888(gWhite);
    }
    return out;
}

1.5.3 棕褐色效果

棕褐色效果的代码与单色效果的代码非常相似。唯一的区别在于像素的算法计算。使用不同的公式和常数来获得最终的像素数据。以下是棕褐色效果运行 OpenCL 和原生 C 代码的函数声明。如您所见,除了名称差异外,函数声明和定义是相同的。

extern int executeSepiaOpenCL(cl_uchar4 *srcImage, cl_uchar4 *dstImage, it int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight);

extern int executeSepiaNative(cl_uchar4 *srcImage, cl_uchar4 *dstImage, int radiHi, int radiLo, int xt, int yt, int nWidth, int nHeight);

JNIEXPORT jstring JNICALL Java_com_example_imageprocessingoffload_NativeLib_renderSepia(JNIEnv * env, jclass obj, jobject bitmapIn, jobject bitmapOut, jint renderocl, jlong time_ms, jstring ename, jint xto, jint yto, jint radHi, jint radLo, jint devtype, jobject assetManager) { … }

Java_com_example_imageprocessingoffload_NativeLib_renderSepia(…) 中的源代码片段与单色示例非常相似,因此省略。

当调用 executeSepiaOpenCL(…) 时,该函数将 pixelsInpixelsOut 强制转换为 cl_uchar4 缓冲区并传递。该函数使用 OpenCL API 创建缓冲区和其他适当的资源。它设置内核参数并排队必要的命令来执行 OpenCL 内核。由 pixelsIn 指向的图像输入缓冲区被分配为只读缓冲区。内核代码使用 pixelsIn 缓冲区指针获取像素数据。像素数据被内核算法用于将传入图像转换为单色图像。输出缓冲区是读写缓冲区,它保存图像结果并由 pixelsOut 指向。

executeSepiaNative(…) 函数包含用 C 代码编写的棕褐色算法。该算法是基本的,由一个外循环(y 循环)和一个内循环(x 循环)组成,用于计算像素数据,其结果存储在由 pixelsOut 指向的 dstImage 中。由 pixlesIn 指向的 srcImage 用于解引用输入像素数据,以便算法公式将其转换为单色像素。

棕褐色效果的 OpenCL 内核代码

constant uchar4 cWhite = {1, 1, 1, 1};
constant float3 sepiaRed = {0.393f, 0.769f, 0.189f};
constant float3 sepiaGreen = {0.349f, 0.686f, 0.168f};
constant float3 sepiaBlue = {0.272f, 0.534f, 0.131f};

__kernel void sepia(__global uchar4 *in, __global uchar4 *out, int4 intArgs, int2 wh)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int width = wh.x;
    int height = wh.y;
    
    if(width <= x || height <= y) return;
    
    int xTouchApply = intArgs.x;
    int yTouchApply = intArgs.y;
    int radiusHi = intArgs.z;
    int radiusLo = intArgs.w;
    int tid = x + y * width;
	
    uchar4 c4 = in[tid];
    float4 f4 = convert_float4(c4);
    int xRel = x - xTouchApply;
    int yRel = y - yTouchApply;
    int polar = xRel*xRel + yRel*yRel;
    
    uchar4 pixOut;
      
    if(polar > radiusHi || polar < radiusLo)
    {
        if(polar < radiusLo)
        {
        	float4 outPixel;
            float tmpR = dot(f4.xyz, sepiaRed);
            float tmpG = dot(f4.xyz, sepiaGreen);
            float tmpB = dot(f4.xyz, sepiaBlue);
            
            outPixel = (float4)(tmpR, tmpG, tmpB, f4.w);
            pixOut = convert_uchar4_sat_rte(outPixel);
        }
        else
        {
            pixOut= c4; 
        }
    }
    else
    {
         pixOut = cWhite; 
    }
    out[tid] = pixOut;
}

The RenderScript kernel code for the sepia effect:

#pragma version(1)
#pragma rs java_package_name(com.example.imageprocessingoffload)
#pragma rs_fp_relaxed

int radiusHi;
int radiusLo;
int xTouchApply;
int yTouchApply;

rs_script gScript;
const float4 gWhite = {1.f, 1.f, 1.f, 1.f};

const static float3 sepiaRed = {0.393f, 0.769f, 0.189f};
const static float3 sepiaGreen = {0.349f, 0.686, 0.168f};
const static float3 sepiaBlue = {0.272f, 0.534f, 0.131f};

uchar4 __attribute__((kernel)) sepia(uchar4 in, uint32_t x, uint32_t y)
{
    uchar4 result;
    float4 f4 = rsUnpackColor8888(in);
    
    int xRel = x - xTouchApply;
    int yRel = y - yTouchApply;
    int polar = xRel*xRel + yRel*yRel;
    
    if(polar > radiusHi || polar < radiusLo)
    {
    	if(polar < radiusLo)
      {
        	float3 out;
        	       	
        	float tmpR = dot(f4.rgb, sepiaRed);
        	float tmpG = dot(f4.rgb, sepiaGreen);
        	float tmpB = dot(f4.rgb, sepiaBlue);
        	
        	out.r = tmpR;
        	out.g = tmpG;
        	out.b = tmpB;
        	result = rsPackColorTo8888(out);
        }
        else 
        {
            result = rsPackColorTo8888(f4);
        }
    }
    else  
    {
         result = rsPackColorTo8888(gWhite);
    }
    return result;
}

1.6 运行 RenderScript 的代码和资源

RenderScript 实现需要什么才能执行图像效果?虽然不是规则甚至建议,但为了简单起见,示例应用程序使用了在全局范围内定义的常见资源和变量。Android 开发者可以根据应用程序的复杂性使用不同的方法来定义常见资源。

以下常用资源和全局变量在 MainActivity.java 源文件中声明和定义。

private RenderScript rsContext;

rsContext 变量对所有脚本通用,用于存储 RenderScript 上下文。上下文作为 RenderScript 框架的一部分进行设置。要了解更多内部工作原理,请参阅 RenderScript 框架

private ScriptC_plasma plasmaScript;
private ScriptC_mono monoScript;
private ScriptC_sepia sepiaScript;
	
The plasmaScript, monoScript, and sepiaScript variables are instances of the class that wraps access to the specific RenderScript kernels. Eclipse* IDE automatically generates the Java class from the rs files i.e., ScriptC_plasma from plasma.rs, ScriptC_mono from mono.rs, and ScriptC_sepia from sepia.rs. Specific RenderScript wrapper classes are generated and can be located in Java files under the gen folder. For example, for the sepia.rs file, the Java class is found in the ScriptC_sepia.java file. To generate the Java code, the rs file must completely define the RenderScript kernel code and be syntactically correct to compile. For the sample application, all ScriptC_<*> classes were imported in the MainActivity.java code.
private Allocation allocationIn;
private Allocation allocationOut;
private Allocation allocationPalette;
private Allocation allocationAngles;

分配是 RenderScript 内核操作的内存抽象。例如,allocationInallocationOut 存储输入和输出图像的纹理数据。示例应用程序中脚本的输入是 AllocationIn,而 AllocationOut 是保存由 RenderScript 内核生成的图像数据的输出。Palette 和 Angles 分配用于将角度和查找表数据传递给内核。数据是在调用等离子效果的 RenderScript 之前在主活动代码中生成的。PaletteAngles 数据是生成等离子效果媒体所必需的。

用于将资源和生成代码粘合在一起以运行 RenderScript 内核的代码在示例应用程序的 initRS(…) 辅助函数中定义。

protected void initRS() { … };

initRS() 函数通过 RenderScript 对象的 create 方法初始化 RenderScript 上下文。如前所述,上下文句柄对所有渲染脚本通用,并存储在 rsContext 全局变量中。RenderScript 对象的实例化需要 RenderScript 上下文。以下代码行在示例应用程序 MainActivity 的范围内创建 RenderScript 上下文,因此“this”被传递给 RenderScript.create(…) 方法调用。

rsContext = RenderScript.create(this);

创建 RenderScript 上下文后,将分配执行内核代码所需的特定应用程序 RenderScript 对象。以下源代码行显示了 initRS() 函数中的逻辑,该函数根据需要实例化 RenderScript 对象。

if (effectName.equals("plasma")) {
plasmaScript = new ScriptC_plasma(rsContext);
} else if (effectName.equals("mono")) {
	monoScript = new ScriptC_mono(rsContext);
} else if (effectName.equals("sepia")) {
	sepiaScript = new ScriptC_sepia(rsContext);
} // add here to add additional effects to the application

stepRenderScript(…) 是一个辅助函数,用于运行给定效果的 RenderScript。它使用 RenderScript 对象设置所需参数并调用 RenderScript 内核。下面的源代码是 stepRendeScript(…) 函数的一部分,显示了如何为等离子和单色效果调用 RenderScript 内核。

private void stepRenderScript(…) {

 … // code omitted for simplification
 if(effectName.equals("plasma")) {
	plasmaScript.bind_gPalette(allocationPalette);
	plasmaScript.bind_gAngles(allocationAngles);
	plasmaScript.set_gx(inX - stepCount);
	plasmaScript.set_gy(inY - stepCount);
	plasmaScript.set_ts(System.currentTimeMillis() - mStartTime);
	plasmaScript.set_gScript(plasmaScript);
	plasmaScript.invoke_filter(plasmaScript, allocationIn, allocationOut);
 }
 else if(effectName.equals("mono")) {
// Compute parameters "circle of effect" depending on number of elapsed steps.
	int radius = (stepApply == -1 ? -1 : 10*(stepCount - stepApply));
	int radiusHi = (radius + 2)*(radius + 2);
	int radiusLo = (radius - 2)*(radius - 2);
	// Setting parameters for the script.
	monoScript.set_radiusHi(radiusHi);
	monoScript.set_radiusLo(radiusLo);
	monoScript.set_xInput(xToApply);
	monoScript.set_yInput(yToApply);
	// Run the script.
	monoScript.forEach_root(allocationIn, allocationOut);
	if(stepCount > FX_COUNT)
	{
		stepCount = 0;
		stepApply = -1;
	}
 }
 else if(effectName.equals("sepia")) {
    … // code similar to mono effect
 }
 … // code omitted for simplification

};

gPalettegAnglesgxgygScript 是在等离子 RenderScript 内核中定义的全局变量。RenderScript 框架生成函数以将所需数据传递给内核运行时。所有变量都在 plasma.rs 文件中声明。定义为 rs_allocation 的变量生成 bind_<var> 函数。对于等离子效果,生成 bind_<gvars> 函数以将 PaletteAngles 数据绑定到 RenderScript 上下文。对于标量参数,例如 gxgytsgScript,生成 set_<var> 方法以发送该参数的特定数据。标量参数用于发送等离子 RenderScript 内核所需的正在运行的 x、y 值和时间戳。invoke_filter(…) 函数是根据 RenderScript 定义生成的。在等离子脚本中定义用户函数(如 filter() 函数)是一种编程可配置和/或可重用 RenderScript 内核代码的方式。

对于单色效果,radius 用于计算 radiusHiradiusLo 参数。这些参数与 xInputyInput 一起用于计算和显示单色效果的径向进度。请注意,对于单色脚本,不是调用用户函数,而是直接调用 forEach_root()forEach_root(…) 是默认方法,由框架为渲染脚本生成。请注意,radiusHiradiusLoxInputyInput 在内核代码中定义为全局变量,并且生成 set_<var> 方法以将所需数据传递给 RenderScript 内核。

如需更多帮助,请参阅 RenderScript 源代码定义。

2. OpenCL 包装类

包装类提供用于 OpenCL API 的函数,以编译和执行 OpenCL 内核。它还为 API 提供包装函数以初始化 OpenCL 运行时。包装类的目的是促进运行时环境的初始化和设置,以执行 OpenCL 内核。以下是包装类中每个方法的简要描述和用法。使用下载链接获取 OpenCL 包装类的完整源代码。

class openclWrapper {
private:
cl_device_id* mDeviceIds;	// Holds OpenCL device Ids (CPU, GPU, etc...)
	cl_kernel mKernel;		// Holds handle for kernel to run
	cl_command_queue mCmdQue;	// Holds command queue for CL device
	cl_context mContext;		// Holds OpenCL context
	cl_program mProgram;		// Holds OpenCL program handle

public:
	openclWrapper() {
		mDeviceIds = NULL;
		mKernel = NULL;
		mCmdQue = NULL;
		mContext = NULL;
		mProgram = NULL;
	};
	~openclWrapper() { };
	cl_context getContext() { return mContext; };
	cl_kernel getKernel() { return mKernel; };
	cl_command_queue getCmdQue() { return mCmdQue; };

	int createContext(cl_device_type deviceType);
	bool LoadInlineSource(char* &sourceCode, const char* eName);
	bool LoadFileSource(char* &sourceCode, const char* eName, AAssetManager *mgr);
	int buildProgram(const char* eName, AAssetManager *mgr);
	int createCmdQueue();
	int createKernel(const char *kname);
	// overloaded function
	int initOpenCL(cl_device_type clDeviceType, const char* eName, AAssetManager *mgr=NULL);
};
  • ::createContext(cl device) 函数 - 是一个辅助函数,它使用设备类型(例如,CPU 或 GPU)来验证 OpenCL 支持并从系统中获取设备 ID。此函数使用设备 ID 来创建 OpenCL 执行上下文。该函数作为 OpenCL 初始化步骤的一部分被调用。它返回 SUCCESS 并设置类上下文句柄,即 mContext,如果平台或设备 ID 枚举和/或上下文本身的创建失败,则返回 FAIL。
  • ::createCmdQue() 函数 - 枚举与 CL 上下文关联的设备数量。它依赖于私有数据成员 mContext 来创建命令队列。如果成功,返回 SUCCESS 并设置命令队列句柄,即 mCmdQue;如果无法为之前由 createContext(…) 函数枚举的设备 ID 创建命令队列,则返回 FAIL。
  • ::buildProgram(effectName, AssetManager) 函数 - 是一个重载函数,它接受图像处理算法名称(也定义为效果名称)和指向 Android JNI 资产管理器的指针。资产管理器使用效果名称定位并读取包含内核源代码的 OpenCL 文件。包装类还使用效果名称定位并加载“内联”OpenCL 源代码。该函数被重载,因为其声明默认将资产管理器指针设置为 NULL。本质上,该函数只能与效果名称一起调用,或者与效果和指向资产管理器的有效指针一起调用,以决定何时编译内联定义的 OpenCL 代码或从单独文件中加载 OpenCL 代码。这允许程序员将 OpenCL 程序定义和部署为内联字符串或在单独的 OpenCL 文件中。资产管理器指针值用于调用从字符串加载 OpenCL 程序的函数,或者调用使用资产管理器 API 将 OpenCL 源代码读入缓冲区的函数。
    • buildProgram(…) 函数调用 OpenCL API clCreateProgramWithSource(…) 以从源代码创建程序。如果 OpenCL 源代码存在语法错误,create program with source API 将返回错误并无法创建程序。OpenCL 上下文和源代码缓冲区作为参数传递。如果 CL 程序编译成功,clCreateProgramWithSource(…) 将返回程序句柄。
    • clBuildProgram(…) API 接受由 clCreateProgramWithSource(…)clCreateProgramWithBinary(…) API 创建的程序句柄。调用 clBuildProgram(…) 来编译和链接将在 CL 设备上运行的可执行程序。如果出现错误,可以使用 clGetProgramBuildInfo(…) 来转储编译错误。有关示例,请参阅包装类源代码。
  • ::createKernel(…) 函数 - 接受效果名称并使用程序对象创建内核。如果内核创建成功,函数返回 SUCCESS。有效的内核句柄存储在 mKernel 中,该句柄随后用于设置内核参数并执行实现图像处理算法的 OpenCL 内核。
  • ::getContext()::getCmdQue()::getKernel() 方法仅返回上下文、命令队列和内核句柄。这些句柄在 JNI 函数中用于将所需命令排队以运行 OpenCL 内核。

3. 总结

本文重点介绍了可用于在 Android 应用程序中卸载图像处理的一些 OpenCL 技术和过程。与 RenderScript 类似,OpenCL 是一种可行且强大的技术,可以卸载您的图像处理工作负载。随着越来越多的设备支持 OpenCL,很高兴知道这种语言技术可以卸载并有望加速您的图像处理工作负载。有关更多信息,请参阅 Intel SDK for OpenCL 文档。

4. 作者简介

Eli Hernandez 是英特尔公司消费类客户端和功耗使能部门的应用工程师,他与客户合作优化其软件以提高能效,并在英特尔硬件和软件技术上实现最佳运行。Eli 于 2007 年 8 月加入英特尔,在电信和化学行业的软件开发方面拥有超过 12 年的经验。他于 1989 年获得电气工程学士学位,并于 1991-1992 年在芝加哥德保罗大学完成计算机科学硕士研究。

© . All rights reserved.