![](https://habrastorage.org/files/38f/f10/b7c/38ff10b7c857489ab5d1cb2f70a22482.jpg)
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) помещаются в буфер. Затем эти изображения накладываются одно на другое на основе значения альфа (это фактор веса в уравнении вычисления пикселей). Значение альфа передается в виде параметра.
![](https://habrastorage.org/files/563/f35/e66/563f35e66f1b4fc6a53db52e0fdc1103.png)
Рисунок 1. Альфа = 0,84089642
Входные изображения должны быть 24- или 32-разрядными. На выходе получается 24-разрядное изображение. Входные изображения должны быть одинакового размера. Изображения были в формате ARGB, это учитывалось при их загрузке.
![](https://habrastorage.org/files/455/e2c/6b0/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 в выходном изображении.
Для изменения яркости каждого пикселя можно использовать значение гамма. По умолчанию это значение равно нулю. Пользователь может изменить яркость готового изображения целиком.
Пример запуска программы
![](https://habrastorage.org/files/13e/019/84e/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 невозможна синхронизация зависимостей данных между рабочими группами.
stalkerg
Когда в Beignet будет OpenCL 2.0?