Использование изображений, доступных для чтения и записи, в OpenCL™ 2.0

Благодарности

В редактировании этой статьи и приведенного в ней примера кода нам помогли Хавьер Мартинес, Кевин Пател и Теджас Будух.

Введение

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

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, выходные данные могут быть скопированы во входной буфер. Это позволяет снизить количество необходимых ресурсов.

В OpenCL 1.2 изображения доступны только для чтения или записи. Для каких-либо изменений изображения необходимо обрабатывать изображение как буфер и работать с этим буфером (см. cl_khr_image2d_from_buffer: https://software.intel.com/ru-ru/articles/using-image2d-from-buffer-extension).

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

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

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

Using Alpha value 0.84089642

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

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

Using Alpha value of 0.32453

Рисунок 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 в выходном изображении.

Для изменения яркости каждого пикселя можно использовать значение гамма. По умолчанию это значение равно нулю. Пользователь может изменить яркость готового изображения целиком. 

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

Read Write Image Sample Program running on OCL2.0 Device

Рисунок 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 невозможна синхронизация зависимостей данных между рабочими группами.

Справочные материалы

Об авторах

Олудемилад Раджи (Oludemilade Raji) — инженер по графическим драйверам в отделе Visual and Parallel Computing Group корпорации Intel. Он уже в течение 4 лет занимается языком программирования OpenCL и принимал участие в разработке драйверов для Intel HD Graphics, включая разработку OpenCL 2.0.

 

Роберт Иоффе (Robert Ioffe) работает на должности технического инженера-консультанта в отделе Software and Solutions Group корпорации Intel. Роберт — опытный специалист по программированию OpenCL и оптимизации нагрузки OpenCL на ГП Intel Iris и Intel Iris Pro, он прекрасно знает графическое оборудование Intel. Роберт принимал активное участие в разработке стандартов Khronos, занимался прототипированием самых последних функций и обеспечением их работоспособности в архитектуре Intel. В последнее время Роберт работал над прототипированием возможностей вложенной параллельной обработки (функции enqueue_kernel) OpenCL 2.0, он написал ряд примеров, демонстрирующих эту функциональность, включая алгоритм GPU-Quicksort для OpenCL 2.0. Кроме того, он записал и выпустил два видеоролика по оптимизации ядер OpenCL, сейчас он работает над третьим видеороликом, посвященным возможностям вложенной параллельной обработки.

Para obter informações mais completas sobre otimizações do compilador, consulte nosso aviso de otimização.