[Перевод] Использование изображений, доступных для чтения и записи, в OpenCL 2.015.10.2015 14:33

38ff10b7c857489ab5d1cb2f70a22482.jpgДо OpenCL 2.0 было невозможно проводить операции чтения и записи изображения в рамках одного и того же ядра. Можно было объявлять изображения как CL_MEM_READ_WRITE, но после передачи изображения ядру приходилось выбирать одно из двух: либо __read_only (доступ только для чтения), либо __write_only (доступ только для записи). В OpenCL 2.0 появилась возможность читать и записывать изображения в пределах одного ядра. Однако, имеется несколько особенностей, о которых мы подробно поговорим в этом посте.

input1 = clCreateImage(
oclobjects.context,
CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
&format,
&desc,
&input_data1[0],
&err );
SAMPLE_CHECK_ERRORS( err );

Фрагмент кода 1. Можно было создать буфер изображения с помощью CL_MEM_READ_WRITE

__kernel void Alpha( __read_write image2d_t inputImage1, 
__read_only image2d_t 
inputImage2, 
uint width, 
uint height, 
float alpha, 
float beta, 
int gamma )

Фрагмент кода 2. В OpenCL 2.0 появилась возможность читать и записывать изображения в пределах одного ядра

Преимущества изображений, доступных для чтения и записи


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

В OpenCL 1.2 и более ранних версиях изображения могли иметь только квалификаторы __read_only и __write_only. В OpenCL 2.0 появился квалификатор __read_write, выходные данные могут быть скопированы во входной буфер. Это позволяет снизить количество необходимых ресурсов. Для каких-либо изменений изображения необходимо обрабатывать изображение как буфер и работать с этим буфером (см. cl_khr_image2d_from_buffer).

Текущее решение состоит в том, чтобы обрабатывать изображения как буферы и управлять буферами. Для обработки двухмерных изображений как буферов требуется затратить определенное количество ресурсов. При этом также становится невозможно использовать возможности срезания и фильтрации, доступные в read_images. Поэтому желательно использовать изображения с квалификатором read_write.

Обзор примера


В примере два растровых изображения (input1.bmp и input2.bmp) помещаются в буфер. Затем эти изображения накладываются одно на другое на основе значения альфа (это фактор веса в уравнении вычисления пикселей). Значение альфа передается в виде параметра.

563f35e66f1b4fc6a53db52e0fdc1103.png
Рисунок 1. Альфа = 0,84089642

Входные изображения должны быть 24- или 32-разрядными. На выходе получается 24-разрядное изображение. Входные изображения должны быть одинакового размера. Изображения были в формате ARGB, это учитывалось при их загрузке.

455e2c6b0b674c74bc2fa2075f16c5f6.png
Рисунок 2. Альфа = 0,32453

Формат ARGB преобразуется в RGBA. Изменение значения бета приводит к значительным изменениям выходного изображения.
Использование SDK
В SDK демонстрируется наложение изображений при помощи чтения и записи. Для управления работой образца кода можно использовать следующие параметры командной строки.

Параметры Описание
-h, --help Отображение этого текста и выход.
-p, --platform <число или строка> Выбор платформы, устройства которой используются.
-t, --type all | cpu | gpu | acc | default | <константа OpenCL для типа устройства> Выбор типа устройства, на котором выполняется ядро OpenCL.
-d, --device <число или строка> Выбор устройства, на котором выполняется вся работа.
-i, --infile <24- или 32-разрядный входной BMP-файл> Имя первого прочитываемого файла в формате BMP. По умолчанию — input1.bmp.
-j, --infile <24- или 32-разрядный входной BMP-файл> Имя второго прочитываемого файла в формате BMP. По умолчанию — input2.bmp.
-o, --outfile <24- или 32-разрядный входной BMP-файл> Имя выходного файла, в который производится запись. По умолчанию — output.bmp для OCL1.2 и 20_output.bmp для OCL2.0.
-a, --alpha <значение с плавающей запятой от нуля до единицы> Ненулевое положительное значение, определяющее, насколько два изображения будут наложены одно на другое при совмещении. Значение альфа по умолчанию — 0,84089642. Значение бета по умолчанию — 0,15950358.


В образце SDK заданы значения по умолчанию, благодаря чему приложение может работать без какого-либо ввода со стороны пользователя. Пользователи могут использовать собственные входные BMP-файлы. Все файлы должны быть 24- или 32-разрядными. Значение альфа определяет, насколько первое изображение будет накладываться на второе.

calculatedPixel = ((currentPixelImage1 * alpha) + (currentPixeImage2 * beta) + gamma);


Значение бета равно разности между единицей и значением альфа.

float beta = 1 – alpha;


Эти два значения определяют «вес» изображений 1 и 2 в выходном изображении.
Для изменения яркости каждого пикселя можно использовать значение гамма. По умолчанию это значение равно нулю. Пользователь может изменить яркость готового изображения целиком.

Пример запуска программы


13e01984efc84f948733a8a970a1d698.png
Рисунок 3. Запуск программы на устройстве OpenCL 2.0
 

Ограничения изображений, доступных для чтения и записи


Ограничители нельзя использовать с изображениями, для которых требуется синхронизация между разными рабочими группами. Для свертки изображений требуется синхронизация всех потоков. Свертка по отношению к изображениям обычно предусматривает математические действия над двумя матрицами и создание в результате третьей матрицы. В примере свертки изображений используется размытие Гаусса. В других примерах используется повышение резкости изображений, обнаружение краев и придание рельефности.

В качестве примера рассмотрим размытие Гаусса. Фильтр Гаусса — это низкопроходный фильтр, удаляющий высокочастотные значения. В результате снижается уровень детализации изображения и получается эффект размытия. Применение размытия Гаусса — то же самое, что преобразование изображения с помощью функции гауссова распределения (ее часто называют маской). Для демонстрации функциональности чтения и записи изображений пришлось применить размытие по горизонтали и по вертикали.

В OpenCL 1.2 это пришлось бы делать в два прохода. Одно ядро использовалось бы только для размытия по горизонтали, а другое — для размытия по вертикали. Результат одного размытия будет использоваться в качестве входных данных для следующего (в зависимости от того, какое размытие было первым).

__kernel void GaussianBlurHorizontalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);
    float4 calculatedPixel = (float4)(0,0,0,0);
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(maskIndex, 0));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}

__kernel void GaussianBlurVerticalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);
    float4 calculatedPixel = (float4)(0,0,0,0);  
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(0, maskIndex));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}

Фрагмент кода 3. Ядро размытия Гаусса в OpenCL 1.2

В OpenCL 2.0 эти два ядра можно объединить в одно. Используйте ограничитель, чтобы принудительно завершать размытие по горизонтали или по вертикали перед началом следующего размытия.

__kernel void GaussianBlurDualPass( __read_only image2d_t inputImage, __read_write image2d_t tempRW, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);   
    float4 calculatedPixel = (float4)(0,0,0,0)
    currentPixel = read_imagef(inputImage, currentPosition);
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, currentPosition + (int2)(maskIndex, 0));      
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(tempRW, currentPosition, calculatedPixel);

    barrier(CLK_GLOBAL_MEM_FENCE);

    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(tempRW, currentPosition + (int2)(0, maskIndex));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}

Фрагмент кода 4. Ядро размытия Гаусса в OpenCL 2.0

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

© Habrahabr.ru