Компания
281,87
рейтинг
30 июля 2015 в 16:20

Разработка → Ускорение обработки изображений в Android перевод

image Центральные процессоры и графические ядра современных устройств, работающих под управлением Android, способны на многое. Например, их вычислительную мощность можно направить на обработку изображений.

Для того чтобы это сделать, стоит обратить внимание на технологии OpenCL и RenderScript.

В этом материале рассмотрен пример Android-приложения, в котором показаны методики высокопроизводительной обработки изображений с использованием языков программирования OpenCL и RenderScript. Эти технологии разработаны с прицелом на возможности графического аппаратного обеспечения, рассчитанного на параллельную обработку данных (шейдерных блоков). Они позволяют ускорить работу со значительными объёмами данных и решение задач, предусматривающих большое число повторов команд. Хотя, для ускорения обработки графики в Android-приложениях, вы можете воспользоваться другими технологиями, в этом материале рассматриваются примеры построения инфраструктуры приложения и реализации графических алгоритмов на OpenCL и RenderScript. Здесь так же рассмотрен класс-обёртка для OpenCL API, который позволяет упростить создание и исполнение приложений, работающих с графикой и использующих OpenCL. Использование исходного кода этого класса в ваших проектах не требует лицензирования.

При подготовке этого материала предполагалось, что его читатели знакомы с технологиями OpenCL и RenderScript, владеют приёмами программирования для платформы Android. Поэтому основное внимание мы уделим рассмотрению механизмов ускорения обработки или программного создания изображений.

Для того чтобы увидеть примеры в деле, вам понадобится Android-устройство, настроенное так, чтобы оно могло исполнять OpenCL-код. Ниже мы поговорим и о том, как организовать рабочую среду для OpenCL-разработки, используя Intel INDE и Android Studio.

Обратите внимание на то, что цель этой статьи – показ особенностей кода OpenCL и RenderScript, о других технологиях мы здесь не говорим. Кроме того, запланирован материал, посвящённый анализу производительности приложений, использующих код на OpenCL и RenderScript, исполняющийся на видеочипах (GPU).

1.1 Интерфейс приложения


На экране рассматриваемого приложения присутствует три переключателя, с помощью которых можно выбирать между подсистемами работы с изображениями, использующими RenderScript, OpenCL или машинный код Android. Меню позволяет переключаться между исполнением OpenCL-кода на CPU(центральном процессоре) или на GPU (графическом ядре). Кроме того, из меню можно выбрать графический эффект. Выбор целевого устройства доступен только для OpenCL-кода. Платформа Intel x86 поддерживает выполнение OpenCL и на CPU, и на GPU.

Ниже можно видеть главный экран приложения, на котором отображается эффект плазмы, который генерируется средствами OpenCL.


Главное окно программы, выбор целевого устройства для исполнения OpenCL-кода

В правом верхнем углу окна выводятся показатели производительности. Они отображаются для всех трёх поддерживаемых программой способов работы с графикой.

Показатели производительности включают в себя количество кадров в секунду (FPS), время рендеринга кадра (frame render) и время, необходимое для обсчёта эффекта (effect compute elapsed time).


Показатели производительности

Обратите внимание на то, что это – лишь один из примеров производительности. Показатели зависят от устройства, на которой запускается код.

1.2. Используемые API и SDK


В дополнение к ADT (Android Development Tool, средствам разработки Android, которые включают Android SDK), при разработке примера использовались SDK RenderScript и Intel SDK для OpenCL, рассчитанные на работу в среде Android.

Intel OpenCL SDK основан на спецификации OpenCL и придерживается её положений. Эта спецификация представляет собой открытый, бесплатный для использования стандарт кросс-платформенной разработки. Подробности – на сайте Khronos.

RenderScript появился в ADT 2.2. (API Level 8). Это – платформа для выполнения высокопроизводительных вычислений в среде Android. RenderScript, в основном, рассчитан на выполнение задач, допускающих параллельное выполнение вычислений, однако, пользу из него можно извлечь и на вычислениях, которые выполняются последовательно. Здесь вы можете узнать подробности о RenderScript.

Свежая версия ADT, доступная из открытого репозитория Google, включает в себя пакеты, которые нужно импортировать для использования RenderScript, JNI (Java Native Interface, Интерфейс Java с машинным кодом) и набора API времени выполнения.

Дополнительные сведения о разработке с использованием RenderScript можно найти здесь, материалы по OpenCL – здесь.

1.3 Код вспомогательной инфраструктуры приложения


Инфраструктура рассматриваемого приложения состоит из главной Activity и вспомогательных функций. Здесь мы рассмотрим эти функции и код, служащий для настройки пользовательского интерфейса, выбора графического эффекта, а для OpenCL – и для выбора вычислительного устройства, которое следует использовать.

Здесь рассмотрены две основные вспомогательные функции.

Первая, backgroundThread(), запускает отдельный поток выполнения, из которого периодически вызывается функция, выполняющая пошаговое применение графического эффекта. Эта функция взята из приложения, которое рассматривается в материале Начало работы с RenderScript, подробности об этом примере можно найти здесь.

Вторая функция, processStep(), вызывается из backgroundThread(). Она, в свою очередь, вызывает команды для обработки изображений. Функция полагается на состояние переключателей для определения того, какую именно реализацию алгоритмов следует использовать. В функции processStep(), в зависимости от состояния настроек, вызываются методы для обработки изображений с использованием OpenCL, RenderScript или обычного машинного кода, написанного на C/C++. Так как этот код выполняется в фоновом потоке, пользовательский интерфейс не блокируется, поэтому переключаться между реализациями графических алгоритмов можно в любое время. Когда пользователь выбирает эффект, приложение тут же на него переключается.

// Метод processStep() выполняется в отдельном (фоновом) потоке.
private void processStep() {
	try {
		switch (this.type.getCheckedRadioButtonId()) {
		case R.id.type_renderN:
			oclFlag = 0; // OpenCL выключен
			stepRenderNative();
			break;
		case R.id.type_renderOCL:
			oclFlag = 1; // OpenCL включен
			stepRenderOpenCL();
			break;
		case R.id.type_renderRS:
		      oclFlag = 0; // OpenCL выключен
			stepRenderScript();
			break;
		default:
			return;
		}
	} catch (RuntimeException ex) {
		// Обработка исключения и логирование ошибки
		Log.wtf("Android Image Processing", "render failed", ex);
	}
}

1.4 Определение функций, реализованных в машинном коде, в Java


Рассматриваемое приложение реализует класс NativeLib, в котором определены функции, используемые для вызова с помощью JNI команд уровня машинного кода, реализующих графические эффекты. В приложении показано три эффекта: эффект плазмы (plasma), тонирование изображения в сепию (sepia) и обесцвечивание (monochrome). Соответственно, в классе определены функции renderPlasma(…), renderSepia(…) и renderMonoChrome(…). Эти Java-функции играют роль JNI-точек входа, посредством которых вызывается либо функционал, реализованный в машинном коде, либо – OpenCL-версия графических алгоритмов.

Соответствующая JNI-функция, при запуске выбранного графического эффекта, либо исполняет код, написанный на C/C++, либо настраивает и исполняет программу на OpenCL. Описываемый класс использует пакеты android.graphics.Bitmap и android.content.res.AssetManager. Объекты BitMap используются для отправки графических данных в подсистему их обработки и для получения результатов. Приложение пользуется объектом класса AssetManager для того, чтобы получить доступ к OpenCL-файлам (sepia.cl, например). В этих файлах описаны OpenCL-ядра (kernels), представляющие собой функции, реализующие графические алгоритмы.

Ниже показан код класса NativeLib. Его можно легко расширить для добавления дополнительных графических эффектов, на это указывает комментарий //TODO.

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

public class NativeLib
{
    // Реализовано в 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>(…);
    
    //Загрузка библиотеки
    static {
        System.loadLibrary("imageeffectsoffloading");
    }
}

Обратите внимание на то, что Android-объекты AssetManager и BitMap передаются машинному коду в качестве входных и выходных параметров. Объект AssetManager используется машинным кодом для доступа к CL-файлам, в которых описаны OpenCL-ядра. Объект BitMap содержит пиксельные данные, которые обрабатываются в машинном коде. Тот же тип данных используется и для возврата результата обработки.

Параметр пользовательского интерфейса deviceType применяется для того, чтобы указать целевое устройство, на котором будет исполняться OpenCL-код – CPU или GPU. Для достижения такой гибкости ОС Android должна быть соответствующим образом настроена. Современные процессоры Intel Atom и Intel Core могут исполнять OpenCL-инструкции и самостоятельно, и используя встроенные в систему графические чипы.

Параметр eName задаёт то, какое именно OpenCL-ядро следует скомпилировать и запустить. В приложении каждому графическому эффекту соответствует собственная JNI-функция, в результате передача имени ядра может показаться ненужной. Однако, в одном CL-файле (то же касается и JNI-функции) можно закодировать несколько схожих графических алгоритмов. В подобной ситуации параметр eName мог бы быть использован для указания того, какую конкретную CL-программу (или ядро) нужно компилировать и загружать.

Параметр renderocl играет роль флага, указывающего на то, нужно ли исполнять код OpenCL или машинный код, написанный на C/C++. Его значение интерпретируется как истинное, если пользователь активировал переключатель OpenCL, в противном случае флаг остаётся не установленным.

Параметр time_ms используется для передачи отметки времени (в миллисекундах), которая применяется для вычисления показателей производительности. Эффект плазмы, кроме того, ориентируется на это значение при пошаговом обсчёте изображения.

Другие аргументы отражают особенности реализации графических эффектов, в частности, радиальное расширение обработанной области. Например, параметры simXtouch, simYTouch, radLo и radHi, совместно со значениями ширины и высоты, используются в эффектах тонирования изображения в сепию и обесцвечивания. С их использованием происходит расчёт и показ того, как обработка изображения начинается в некоей точке, после чего область радиально расширяется до тех пор, пока не заполнит всё изображение.

1.5 Определения и ресурсы, необходимые для запуска машинного кода (C или OpenCL)


Здесь мы рассмотрим определения машинных JNI-функций, реализующих эффекты, показанные в примере. Как уже было упомянуто, каждому эффекту соответствует одна функция. Так сделано для того, чтобы не усложнять повествование и более чётко выделить функциональные элементы, применяемые при ускорении обработки изображений с использованием OpenCL.

Здесь есть ссылки на код, написанный на C, сюда же включены фрагменты этого кода. Сделано это в расчёте на сравнение производительности рассматриваемых технологий в будущих версиях примера.

Одной JNI-функции соответствует одна машинная Java-функция. Поэтому очень важно, чтобы JNI-функции были правильно объявлены и определены. В Java SDK есть средство javah, которое помогает генерировать правильные и точные объявления JNI-функций. Это средство рекомендовано использовать для того, чтобы не попадать в сложные ситуации, когда код компилируется правильно, но выдаёт ошибки времени выполнения.

Ниже показаны JNI-функции, соответствующие функциям более низкого уровня для ускоренной обработки изображений. Сигнатуры функций сгенерированы с помощью инструмента javah.

// Объявление сигнатуры новой JNI-функции, являющейся точкой входа в
// соответствующую функцию более низкого уровня
#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

Сигнатуры JNI-функций из файла junk.h затем были добавлены в файл imageeffects.cpp, который реализует подготовку и запуск OpenCL или C-кода. Далее, мы выделяем ресурсы, необходимые при запуске OpenCL-кода или машинного кода для эффектов плазмы, обесцвечивания, тонирования в сепию.

1.5.1 Эффект плазмы


Функция Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(…) является входной точкой для исполнения OpenCL-кода или машинного кода, которые реализуют эффект плазмы. Функции startPlasmaOpenCL(…), runPlasmaOpenCL(…), и runPlasmaNative(…) являются внешними по отношению к коду из файла imageeffects.cpp, они объявлены в отдельном файле plasmaEffect.cpp. Исходный код plasmaEffect.cpp можно найти здесь.

Функция renderPlasma(…), являющаяся точкой входа, использует OpenCL-обёртку для запроса у Android поддержки OpenCL. Она вызывает функцию класса-обёртки ::initOpenCL(…) для инициализации OpenCL-окружения. В качестве типа устройства, при создании OpenCL-контекста, передаётся CPU или GPU. Менеджер ресурсов Android использует параметр ceName для идентификации, загрузки и компиляции CL-файла с необходимым ядром.

Если удаётся успешно настроить окружение OpenCL, следующим шагом функции renderPlasma(…) является вызов функции startPlasmaOpenCL(), которая выделяет OpenCL-ресурсы и начинает исполнение ядра, реализующего эффект плазмы. Обратите внимание на то, что gOCL – это глобальная переменная, которая хранит экземпляр класса-обёртки OpenCL. Эта переменная видима для всех JNI-функций, являющихся точками входа. Благодаря такому подходу, OpenCL-окружение можно инициализировать при обращении к любому из поддерживаемых графических эффектов.

При демонстрации эффекта плазмы не используются готовые изображения. Всё, что отображается на экране, сгенерировано программным путём. Параметр bitmapIn – это объект класса BitMap, который хранит графические данные, сгенерированные в ходе работы алгоритма. Параметр pixels, передаваемый в функцию startPlasma(…), отображается на растровую текстуру и используется машинным кодом или кодом ядра OpenCL для чтения и записи пиксельных данных, которые выводятся на экране. Еще раз отметим, что объект assetManager используется для доступа к CL-файлу, который содержит OpenCL-ядро, реализующее эффект плазмы.

JNIEXPORT void Java_com_example_imageprocessingoffload_NativeLib_renderPlasma(JNIEnv * env, jclass, jobject bitmapIn, jint renderocl, jlong time_ms, jstring ename, jint devtype, jobject assetManager) { 
… // опущено для упрощения примера	
    
    // блокируется память для BitMapIn и устанавливается указатель “pixels”, который передаётся OpenCL-функции или функции на машинном коде.

    ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixels); 
 
… // опущено для упрощения примера

  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);
… // опущено для упрощения примера
}

Внешняя функция startPlasmaOpenCL(…) создаёт и заполняет буферы Palette и Angles, которые содержат данные, необходимые для создания эффекта плазмы. Для запуска ядра 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 используется как логический, его значение указывает на то, нужно ли генерировать наборы данных Palette и Angles, которые необходимы для работы алгоритма. Код 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))

/* Цветовая палитра, используемая для рендеринга плазмы */
#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);
}

1.5.2 Обесцвечивание изображений


Функция Java_com_example_imageprocessingoffload_NativeLib_renderMonochrome(…) является точкой входа для вызова функций обесцвечивания изображения, реализованных в машинном коде или средствами OpenCL. Функции executeMonochromeOpenCL(…) и executeMonochromeNative(…) являются внешними по отношению к коду из imageeffects.cpp, они объявлены в отдельном файле. Как и в случае с эффектом плазмы, функция, играющая роль точки входа, использует OpenCL-обёртку для выполнения запросов, касающихся поддержки OpenCL, к подсистеме управления устройствами Android. Она же вызывает функцию ::initOpenCL(…), выполняющую инициализацию OpenCL-окружения.

В следующей паре строк кода показано, что функции executeMonochromeOpenCL(…) и executeMonochromeNative(…) объявлены с ключевым словом extern. Это делает их видимыми для NDK-компилятора. Это необходимо, так как данные функции объявлены в отдельном файле.

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);

В отличие от эффекта плазмы, здесь используются входные и выходные изображения. И bitmapIn, и bitmapOut представляют собой растровые изображения в формате ARGB_888. Оба они отображаются на CL-буферы векторов типа cl_uchar4. Обратите внимание на то, что здесь производится приведение типов pixelsIn и pixelsOut, это необходимо для того, чтобы 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)  {

  … // опущено для упрощения примера

   // блокируется память для BitMapIn и устанавливается указатель “pixelsIn”, который передаётся OpenCL-функции или функции на машинном коде.
   ret = AndroidBitmap_lockPixels(env, bitmapIn, &pixelsIn); 

   // блокируется память для BitMapOut и устанавливается указатель “pixelsOut”, который передаётся OpenCL-функции или функции на машинном коде 
   ret = AndroidBitmap_lockPixels(env, bitmapOut, &pixelsOut); 

 … // опущено для упрощения примера
 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);
    // окончание блока, выполняемого, если OCL инициализирован
else
   executeMonochromeNative((cl_uchar4*) pixelsIn,(cl_uchar4*) pixelsOut, radiHi, radiLo, xt, yt, infoIn.width, infoIn.height);
// Окончание блока OpenCL
… // опущено для упрощения примера

}

Когда вызывается функция executeMonochromeOpenCL(…), pixelsIn и pixelsOut преобразуются к типу cl_uchar4-буферов и передаются. Эта функция использует API OpenCL для создания буферов и других нужных для работы ресурсов. Она задаёт аргументы ядра и ставит в очередь необходимые для исполнения OpenCL-ядра команды. Входной буфер изображения формируется как буфер только для чтения (read_only), для доступа к нему используется указатель pixelsIn. Код ядра использует этот указатель для того, чтобы получить входные пиксельные данные изображения. Эти данные, в свою очередь, обрабатываются ядром, входное изображение обесцвечивается. Выходной буфер – это буфер, предназначенный и для чтения, и для записи (read_write), он хранит результаты обработки изображения, на него указывает pixelsOut. Дополнительные сведения об OpenCL можно найти в Руководстве Intel по программированию и оптимизации.

В функции executeMonochromeNative(…) алгоритм обесцвечивания изображений реализован на C. Это очень простой алгоритм, использующий вложенные циклы – внешний (y) и внутренний (x), в которых обрабатываются пиксельные данные, а результат сохраняется в переменной dstImage, на которую указывает pixelsOut. Переменная srcImage, на которую указывает pixelsIn, используется для получения входных пиксельных данных в формуле алгоритма, где цветное изображение и преобразуется в монохромное.

Вот код ядра 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);
    }
}

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(…), она преобразует переданные ей значения к нужному типу и передаёт pixelsIn и pixelsOut в формате cl_uchar4-буферов. Она использует API OpenCL для создания буферов и других необходимых ресурсов. Она же задаёт аргументы для OpenCL-ядра, ставит в очередь команды для исполнения. Входной буфер изображения формируется как буфер только для чтения (read_only), для доступа к нему используется указатель pixelsIn. Код ядра использует указатель для того, чтобы получить пиксельные данные изображения. Эти данные, в свою очередь, обрабатываются ядром, входное изображение окрашивается в сепию. Выходной буфер – это буфер, предназначенный и для чтения, и для записи (read_write), он хранит результаты обработки изображения, на него указывает pixelsOut.

В функции executeSepiaNative(…) находится реализация алгоритма тонирования в сепию на C. Это простой алгоритм, состоящий из пары вложенных циклов – внешнего (y), и внутреннего (x). В цикле производится обработка данных, результаты сохраняются в переменной dstImage, на которую указывает pixelsOut. Переменная srcImage, на которую указывает pixelsIn, используется для получения входных пиксельных данных в формуле алгоритма, где цветное изображение и окрашивается в сепию.

Ниже показан код 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;
}

1.6 Программный код и ресурсы, необходимые для RenderScript


Что нужно для того, чтобы заработала RenderScript-реализация графических алгоритмов? Нельзя сказать, что это – непреложное правило или даже рекомендация, но в нашем примере, для простоты, использованы обычные глобальные ресурсы и переменные. Android-разработчики могут использовать различные методы для работы с ресурсами, основываясь на сложности приложения.

В файле MainActivity.java объявлены и определены следующие обычно используемые ресурсы.

private RenderScript rsContext;

Переменная rsContext хранит контекст RenderScript, она обычно используется всеми RS-скриптами. Контекст задаётся как часть платформы RenderScript. Здесь можно узнать подробности о внутренних механизмах RenderScript.

private ScriptC_plasma plasmaScript;
private ScriptC_mono monoScript;
private ScriptC_sepia sepiaScript;

Переменные plasmaScript, monoScript, и sepiaScript – это экземпляры класса-обёртки, который предоставляет доступ к RS-ядрам. Eclipse IDE и Android Studio автоматически генерируют соответствующие Java-классы на основе rs-файлов. То есть, на основе файла plasma.rs создаётся класс ScriptC_plasma, на основе mono.rs – класс ScriptC_mono. Файл sepia.rs служит основой для класса ScriptC_sepia. Эти RenderScript-обёртки создаются автоматически, их можно найти в папке gen. Например, обёртку для файла sepia.rs можно найти в файле ScriptC_sepia.java. Для того чтобы система могла автоматически сгенерировать Java-код, в rs-файле должно присутствовать полное, синтаксически верное определение RenderScript-ядра, готовое для компиляции. В приложении-примере классы-обёртки импортированы в код MainActivity.java.

private Allocation allocationIn;
private Allocation allocationOut;
private Allocation allocationPalette;
private Allocation allocationAngles;

Объекты класса Allocation предназначены для организации обмена данными с RenderScript-ядрами. Например, allocationIn и allocationOut хранят графические данные входных и выходных изображений. На вход скрипта подаётся то, на что ссылается allocationIn, а в allocationOut, после обработки в RS-ядре или ядрах, записывается то, что получилось.

Перед запуском RenderScript-ядра, ответственного за вывод эффекта плазмы, в коде главной Activity осуществляется подготовка входных данных. В частности, для создания этого эффекта необходимы данные, хранящиеся в allocationPalette и allocationAngle.

Код, необходимый для запуска RS-ядра, связывает воедино подготовленные ресурсы и автоматически сгенерированные RS-обёртки. Он определён во вспомогательной функции initRS(…).

protected void initRS() { … };

Эта функция инициализирует контекст RenderScript, используя статический метод create класса RenderScript. Как уже было сказано, контекст нужен для работы RenderScript, ссылка на него хранится в глобальной переменной. Контекст RenderScript требуется для создания экземпляров RS-объектов. Ниже показан код, с помощью которого контекст RenderScript создаётся в области видимости MainActivity нашего приложения. Поэтому при вызове метода RenderScript.create(…) передаётся «this».

rsContext = RenderScript.create(this);

После того, как создан RS-контекст, начинается работа с конкретными 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);
} // сюда можно добавить код для реализации дополнительных эффектов

Функция stepRenderScript(…) играет вспомогательную роль, она вызывается для запуска RenderScript-ядра выбранного эффекта. Она использует RenderScript-объекты для установки необходимых параметров и запускает выполнение RS-ядер. Ниже показана часть кода функции stepRenderScript(…), ответственная за настройку и запуск эффектов плазмы и обесцвечивания изображения.

private void stepRenderScript(…) {

 … // опущено для упрощения примера
 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")) {
// Вычисление параметров для радиального применения эффекта, которые зависят от количества уже выполненных шагов 
	int radius = (stepApply == -1 ? -1 : 10*(stepCount - stepApply));
	int radiusHi = (radius + 2)*(radius + 2);
	int radiusLo = (radius - 2)*(radius - 2);
	// Установка параметров скрипта.
	monoScript.set_radiusHi(radiusHi);
	monoScript.set_radiusLo(radiusLo);
	monoScript.set_xInput(xToApply);
	monoScript.set_yInput(yToApply);
	// Запуск скрипта.
	monoScript.forEach_root(allocationIn, allocationOut);
	if(stepCount > FX_COUNT)
	{
		stepCount = 0;
		stepApply = -1;
	}
 }
 else if(effectName.equals("sepia")) {
    … // код, аналогичный таковому для эффекта обесцвечивания
 }
 … // опущено для упрощения примера
};

В RenderScript-ядре, ответственном за эффект плазмы, определены глобальные переменные gPalette, gAngles, gx, gy и gScript. Платформа RS генерирует функции, которые позволяют передавать необходимые данные ядру. Все переменные объявлены в файле plasma.rs. На основе переменных, определенных как rs_allocation, генерируются функции вида bind_<var>. В случае с эффектом плазмы, функции вида bind_<gvars> создаются для осуществления привязки данных, хранящихся в allocationPalette и allocationAngles к RenderScript-контексту. Для скалярных аргументов, таких, как gx, gy, ts и gScript, создаются методы вида set_<var>, которые позволяют отправлять данные для записи в соответствующие переменные. Скалярные параметры, в нашем случае, применяются для передачи в RenderScript-ядро таких данных, как x, y и отметка времени, нужная для эффекта плазмы. Функция invoke_filter(…) автоматически создаётся на основе определения RenderScript. Объявление пользовательских функций, таких, как filter(), в скрипте, генерирующем эффект плазмы, это способ создания настраиваемого и подходящего для повторного использования кода ядра RenderScript.

В случае с эффектом обесцвечивания, параметр radius используется для вычисления аргументов radiusHi и radiusLo. Эти аргументы, вместе с xInput и yInput, используются для обсчёта и вывода на экран результатов применения эффекта в виде радиально расширяющейся области. Обратите внимание на то, что при работе со скриптом, обесцвечивающим изображение, вместо вызова пользовательской функции, forEach_root() вызывается напрямую. Метод forEach_root(…) –, это стандартный метод, создаваемый платформой RenderScript. Обратите внимание на то, что radiusHi, radiusLo, xInput и yInput объявлены в коде ядра как глобальные переменные. Методы вида set_<var> создаются автоматически для передачи данных в ядро.

Взгляните на исходный код эффектов на RenderScript для того, чтобы лучше разобраться в этом примере.

Код ядра RenderScript для реализации эффекта плазмы:

#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);
}

Код ядра RenderScript для обесцвечивания изображения:

#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;
}

Код RenderScript-ядра для применения к изображению эффекта сепии.

#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.7 Настройка рабочей среды и устройств


Для того чтобы приступить к программированию под OpenCL, вам понадобится настроить рабочую среду и устройство для запуска кода.

Для настройки рабочей среды можно воспользоваться пакетом Intel INDE. В ходе его инсталляции можно выбрать IDE, с которой вы планируете работать. В нашем случае это – Android Studio. Когда установка будет завершена, вы получите, во-первых, IDE, готовую к работе (а так же – Android SDK и NDK, образы эмуляторов), во-вторых – набор OpenCL-библиотек и средства настройки устройств для исполнения OpenCL-приложений. Это может быть Android-эмулятор, планшет или смартфон. Нужно учитывать, что для успешной настройки аппаратного устройства нужны Root-права.

Рассмотрим последовательность шагов, необходимую для запуска примера, описанного в документе Руководство по работе OpenCL на платформе Android. Код в руководстве подготовлен с использованием Eclipse IDE, мы же хотим, чтобы с ним можно было работать в Android Studio.

После импорта в Android Studio исходный проект нуждается в доработке. Набор файлов проекта, который подходит для импорта в Android Studio, можно скачать здесь. Если конфигурация вашей системы отличается от той, на которой работал этот пример, обратите внимание на пути к папкам Android SDK, NDK и к папке, где содержатся инструменты Intel для работы с OpenCL.

В частности, в файле Android.mk должен присутствовать актуальный путь к корневой директории, в которой находятся OpenCL-инструменты. В нашем случае соответствующая строка выглядит так:

INTELOCLSDKROOT="C:\Intel\INDE\code_builder_5.1.0.25"

В файле local.properties должны содержаться актуальные пути к Android SDK и NDK.

sdk.dir=C\:\\Intel\\INDE\\IDEintegration\\android-sdk-windows
ndk.dir=C\:\\Intel\\INDE\\IDEintegration\\android-ndk-r10d

После проверки путей можно приступить к настройке Android-эмулятора. Мы пользовались виртуальным устройством Intel Nexus 7 x86. Его нужно запустить из окна Android Virtual Device Manager.

После того, как эмулятор запущен, нам необходимо выяснить его имя, которое требуется для настройки поддержки OpenCL. Узнать его можно, нажав на кнопку Run в Android Studio и дождавшись появления окна выбора устройства. Нас интересует то, что находится в графе Serial Number. В нашем случае это строка emulator-5554.
Теперь откроем командную строку Windows и выполним следующую команду:

C:\Intel\INDE\code_builder_5.1.0.25\android-preinstall>opencl_android_install –d emulator-5554

Она запускает процесс установки OpenCL-библиотек на устройство, имя которого указано в команде. После того, как система сообщит об успешной установке, можно вернуться к окну выбора устройства Android Studio и, выбрав эмулятор, нажать на кнопку OK. Приложение запустится.


Тестирование OpenCL-приложения при помощи эмулятора

Если эмулятор был закрыт, процесс установки OpenCL-библиотек нужно будет повторить.
Та среда разработки, которую мы установили и настроили для работы с OpenCL, подходит и для RenderScript-разработки. Импортируем в Android Studio Eclipse-проект из Руководства по работе с RenderScript на Android. Он аналогичен тому, который был рассмотрен выше, но реализован средствами RenderScript. Для его запуска дополнительные настройки не требуются.

Здесь можно найти дополнительные материалы по OpenCL-разработке. Подробности о настройке среды разработки и устройств для работы с OpenCL вы можете найти здесь.

2. Класс-обёртка OpenCL


Класс-обёртка предоставляет функциональность OpenCL API для компиляции и исполнения OpenCL-ядер. Кроме того, он предоставляет функции-обёртки для API, предназначенные для инициализации среды времени выполнения OpenCL. Основная задача класса состоит в помощи при инициализации и настройке среды времени выполнения. Ниже приведено описание методов класса и комментарии по их использованию. Исходный код класса можно загрузить здесь.

class openclWrapper {
private:
cl_device_id* mDeviceIds;	// Хранит идентификаторы OpenCL-устройств (CPU, GPU, и так далее)
	cl_kernel mKernel;		// Идентификатор ядра 
	cl_command_queue mCmdQue;	// Очередь команд для CL-устройства
	cl_context mContext;		// Контекст OpenCL 
	cl_program mProgram;		// Идентификатор OpenCL-программы

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);
	// перегруженная функция
	int initOpenCL(cl_device_type clDeviceType, const char* eName, AAssetManager *mgr=NULL);
};

Функция ::createContext(cl device) играет вспомогательную роль. Она, используя выбранное устройство (CPU или GPU), проверяет, поддерживает ли это устройство OpenCL и получает его идентификатор у системы. Получив идентификатор, она создаёт контекст исполнения OpenCL. Функция вызывается в ходе процесса инициализации OpenCL. Она, при успешном завершении, возвращает SUCCESS и устанавливает идентификатор контекста (mContext). В противном случае, если система не вернула идентификатор устройства, либо не удалось создать контекст OpenCL, функция возвращает FAIL.

Функция ::createCmdQue() работает с устройствами, связанными с OpenCL-контекстом. Она использует закрытый член данных mContext для создания очереди команд. При успешном завершении функция возвращает SUCCESS и устанавливает идентификатор очереди команд (mCmdQue). В противном случае, если не удаётся создать очередь команд для идентификатора устройства, ранее возвращённого функцией createContext(…), возвращается FAIL.

Функция ::buildProgram(effectName, AssetManager) перегружена. Она принимает название алгоритма обработки изображения (effectName) и указатель на менеджер ресурсов Android JNI. Менеджер ресурсов использует название алгоритма для обнаружения и считывания OpenCL-файла, который содержит исходный код соответствующего ядра. Класс-обёртка так же использует название эффекта для поиска и загрузки встроенного (inline) OpenCL-кода. При перегрузке функции в её объявлении ссылка на менеджер ресурсов установлена в NULL по умолчанию. Важное следствие этого заключается в том, что её можно вызвать либо только с параметром effectName, либо и с ним, и с рабочим указателем на менеджер ресурсов. После вызова принимается решение о том, нужно ли компилировать OpenCL-код, определенный с использованием механизма встраивания, либо загружать код из отдельного файла. Такой подход позволяет программисту либо строить и разворачивать программу с использованием встроенного кода, в виде строк, либо – в виде отдельных OpenCL-файлов. Указатель на менеджер ресурсов используется либо для вызова функции, которая загружает OpenCL-программу из строки, либо – функции, которая задействует API менеджера ресурсов для чтения исходного кода из OpenCL-файла в буфер.

  • Функция buildProgram(…) вызывает команду OpenCL API clCreateProgramWithSource(…) для создания программы из исходного кода. Это API возвращает сообщения об ошибках в том случае, если в OpenCL-коде есть синтаксические неточности. Программу при таком состоянии дел создать невозможно. Контекст OpenCL и буфер с исходным кодом передаются в виде аргументов. Если компиляция завершилась успешно, clCreateProgramWithSource(…) возвращает идентификатор программы.

  • Команда clBuildProgram(…) принимает идентификатор программы, полученный из clCreateProgramWithSource(…) или из clCreateProgramWithBinary(…). Эту команду вызывают для того, чтобы скомпилировать и скомпоновать исполняемый файл программы, который затем планируется запускать на OpenCL-устройстве. Если на данном этапе что-то пойдёт не так, можно воспользоваться clGetProgramBuildInfo(…) для сохранения данных по ошибкам компиляции. Соответствующий пример можно найти в исходном коде класса-обёртки.

Функция ::createKernel(…) принимает имя эффекта и использует объект программы для создания ядра. В случае успешного создания ядра она возвращает SUCCESS. Идентификатор ядра сохраняется в переменной mKernel, которая затем используется для задания аргументов ядра, которое реализует соответствующий графический алгоритм, и для его исполнения.

Методы ::getContext(), ::getCmdQue() и ::getKernel() возвращают, соответственно, идентификаторы контекста, очереди команд и ядра. Они используются функциями JNI для постановки в очередь команд, необходимых для исполнения OpenCL-ядер.

Итоги


В этом материале освещены некоторые приёмы OpenCL-программирования, которые можно использовать для ускорения обработки графики в Android-приложениях. OpenCL похожа на RenderScript. Это жизнеспособная и мощная технология, способная взять на себя обработку графической информации, которая требует большой вычислительной мощности. С распространением поддержки OpenCL в большем количестве устройств, полезно знать о её возможностях по работе с графикой, по использованию для этих целей не только ресурсов центрального процессора, но и графического ядра. Надеемся, эти знания помогут вам в создании быстрых и качественных приложений.
Автор: @CooperMaster Eliseo Hernandez
Intel
рейтинг 281,87

Комментарии (3)

  • +1
    Было бы интересно сравнить производительность ядер на типах данных buffer и image
  • +1
    Центральные процессоры и графические ядра современных устройств, работающих под управлением Android, способны на многое.

    От Android там только HAL, непосредственно управляет процессорами и периферией в SoC ядро Linux, Андроид тут добавляет Java-прослойку. Впрочем она и определяет «способность на многое» количеством Java-программистов. Тем не менее познавательная информация про интеграцию OpenCL kernels в приложения для Андроид.
  • 0
    Очень подробная инструкция о том, как разрабатывать OpenCL код для Андроид.
    И ни слова о том, что эта технология официально не поддерживается в данной ОС.
    Гугл в качестве альтернативы OpenCL выдвинул свой RenderScript.

Только зарегистрированные пользователи могут оставлять комментарии. Войдите, пожалуйста.

Самое читаемое Разработка