Расчет гистограммы изображения на GPU. Атомарные типы данных Metal и операции над ними

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

Чтобы хоть немного минимизировать издержки, а результат применения CLUT сделать более предсказуемым и воспроизводимым, как правило над изображением производится некий препроцессинг, часто под ним подразумевают нормализацию. Нормализация, к примеру, почти всегда используется в Computer Vision и прочих не имеющих отношения к фотографии приложениях Image Processing в своих корыстных целях. Пока же, мы в основном говорим о Image Processing-е как об автоматизированной обработке фотографий для изменения их перцептуальных свойств. Но что бы их изменить автоматически нам нужно как-то научиться эти свойства анализировать и использовать результаты этого анализа в настройке наших фильтров.

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

Я не буду погружаться в бессмысленное описание деталей связанных с этим инструментом, напомню, что в основном под гистограммой понимают графическое представление распределения интенсивностей  изображения и как правило, если оно цветное, то в пространстве RGB или RGBA или, как мне больше нравится, RGBY, где Y-канал яркости стыренный из пространства YCbCr. А еще мы будем понимать под гистограммой не графическое представление, а некий контейнер данных, над которым возможны операции и извлечение из этих данных свойств за разумное машинное время достаточное для анализа в режиме реального времени. Под самим представлением данных мы будем понимать поканальное вероятностное распределение каких то величин. А под величинами понимать интенсивность.

Получение гистограммы изображения в iOS

Получить гистограмму изображения в iOS (и в OS X, тоже) можно элементарно. Как и в любой другой среде. Первое, что приходит на ум — это просто взять и написать что-то вроде:

...
for(int u=0; u->height; u++){
for(int v=0; v&u->width; v++){
histogram.bins[data[u][v]]++;
}
}
...

А можно использовать Accelerate Framework:

...
vImage_Buffer buffer;
buffer.data = CVPixelBufferGetBaseAddress(pixelBuffer);
buffer.width = CVPixelBufferGetWidth(pixelBuffer);
buffer.height = CVPixelBufferGetHeight(pixelBuffer);
buffer.rowBytes = CVPixelBufferGetBytesPerRow(pixelBuffer);

vImage_Error error = vImageHistogramCalculation_ARGB8888(&buffer, histogram, 0);
...

Можно подключить к проекту: GPUImage.

А можно воспользоваться новым Metal Performance Shaders и собрать гистограмму прямо на GPU:

...
var histogramInfo = MPSHistogramInfo( numberOfHistogramEntries: 256, histogramForAlpha: false, minPixelValue: float4( 0, 0, 0, 0 ), maxPixelValue: float4( 1, 1, 1, 1 ) )

histogram = MPSImageHistogram( device: device, histogramInfo: histogramInfo )

histogramBuffer = device.newBufferWithLength( 256 * 4 * sizeof( Int32 ), options:[] )

histogram.encodeToCommandBuffer(commandBuffer, sourceTexture: sourceTexture, histogram: histogramBuffer!, histogramOffset: 0)

commandBuffer.commit()
...

Достоинства и недостатки различных методов получения гистограммы в iOS

Очевидно, что у первого метода вычисления нет никаких достоинств кроме простоты. Тут даже обсуждать ничего не будем: медленно и неэффективно. Второй метод предоставляет практически молниеносное вычисление гистограммы если данные изображения уже загружены в память CPU.  А ведь это еще нужно провернуть. Поэтому я считаю его устаревшим и в конечном счете неэкономичным. Ну сами посудите: нам нужно что-то сделать с картинкой в GPU, мы её уже загнали в текстуру этого GPU, чтобы произвести какие-то непристойные глумления, и вот теперь мы обязаны вытащить все это еще раз назад в память процессора, обернуть структуру понятную специальному методу специального фреймворка, и потом еще проследить за корректным освобождением памяти. Вычеркиваем.

GPUImage — отличный, распространенный, и хорошо отлаженный тулкит для работы с изображениями. Использует OpenGL ES и стало быть позволяет быстро распараллелить вычисления и собрать гистограмму на GPU. Но наша цель Metal!.

Наконец остается новый framework MPS. Тут нам не нужно будет делать лишних телодвижений — просто исполняем стандартный танец с бубнами вокруг пайплайна Metal и все. Однако и  тут есть один момент — мы никак не управляем деталями расчета. А если захотим отгистрограмить хитрое цветовое пространство или вообще посчитать распределения цветов в HSV круге для вычисления каких-то специальных статистик? Поэтому решим для себя, что пожалуй неплохо бы, повторить это упражнение самостоятельно и собрать гистограмму на GPU самим, так же как это сделали разработчики Apple.

Думаете это просто?

Чорта с два — это очень просто! Хотя до недавнего времени процедура сборки на GPU была не совсем очевидной задачей: статья номер раз, и даже с помощью cuda: статья номер два. Вся сложность тут в том, что распарралелившись по сути по нескольким десяткам или сотням вычислительных потоков нужно уметь сложить результаты в одномерную структуру гистограммы, т.е. придумать механизм устранения конфликтов доступа. В OpenGL таких механизма нет, хотя сами GPU аппаратно его уже давно реализуют. В Cuda есть, но программного слоя для iOS нет. Поэтому нам, в очередной раз повезло и такой механизм таки появился в Metal. Как ни странно называется он атомарные операции. А именно через атомарные типы данных и операции обернутые в Metal Shading Language. MSL предоставляет полноценный интерфейс к ним в стиле C11. Атомарные операции, как известно, могут быть безопасно использованы для доступа к одним и тем же переменным или объектам памяти из разных потоков с сохранением целостности вычислений над ними. И при всем при этом операции эти крайне дешевые с точки зрение использования ресурсов. В случае с MSL эти операции возможны только для типов atomic_int/atomic_uint + почти весь стандартный набор методов над ними. Нам интересно использовать только операцию инкрементации и не соблюдать порядок т.е. нам можно использовать тип памяти memory_order_relaxed. Запишем шейдер с kernel-функцией занимающейся неконфликтным подсчетом гистограммы:

///
/// Контейнер счета интенсивностей
///
typedef struct {
//
// MSL предоставляет полноценный интерфейс к атомарным операцией в стиле C++11.
// Атомарные операции как известно могут быть безопасно использованы для доступа
// к одним итем же переменным или объектам памяти из разных потоков. В случае с MSL
// эти операции возможны только для типов atomic_int/atomic_uint.
//
atomic_uint channel[kIMP_HistogramChannels][kIMP_HistogramSize];
}IMPHistogramBuffer;

///
/// Функция счета.
///
kernel void kernel_impHistogramRGBYCounter(
//
// Исходная текстура интенсивности которой нам нужно посчиать.
//
texture2d inTexture  [[texture(0)]],
texture2d outTexture [[texture(1)]],
//
// Вот собственно структурированный кусок памяти который мы используем для
// подсчета бинов гистограммы. В памяти СPU структура инициализирована
// как 2D массив uint-ов.
// Каждая строка массив значений бинов гистограммы.
//
device IMPHistogramBuffer &out [[ buffer(0) ]],
uint2 gid [[thread_position_in_grid]]
)
{
float4 inColor = inTexture.read(gid);

//
// Максимальный индекс каждого канала в массиве бинов
//
constexpr float3 Im(kIMP_HistogramSize - 1);

//
// Вектор преобразования RGB в яркостный канал Y из YCbCr
//
constexpr float3 Ym(0.299, 0.587, 0.114);

//
// Его тоже растянем до размерности гистограмы
//
uint Y = uint(dot(inColor.rgb,Ym) * inColor.a * Im.x);

//
// Индексы каналов rgb
//
uint3 rgb = uint3(inColor.rgb * Im);

//
// Каждый объект инкремента - указатель на участок памяти в 2D массиве номер строки == номеру канала,
// в нашем конкретном случае RGB,Y.
// Номер индекса - значении бина которое мы будем инкрементировать каждым попавшим в него значением интенсивности
// конкретного канала.
// И вот тут нам очень облегчает жизнь объявление переменной адресов как атомарного типа - такая возможность
// гарантирует неконфликтную инкрементацию значений в множественном массиве вычислительных потоков GPU.
//
// Если бы такой возможности в MSL не существовало, нам бы пришлось попотеть: использовать подход рекурсивной
// обработки текстуры в графических шейдерах, как это, например, описано в:
// http://www.shaderwrangler.com/publications/histogram/histogram_cameraready.pdf
//

//
// Для инкрементации каждой ячейки используем атомарные операции!
//
atomic_fetch_add_explicit(&out.channel[0][rgb.r], 1, memory_order_relaxed);
atomic_fetch_add_explicit(&out.channel[1][rgb.g], 1, memory_order_relaxed);
atomic_fetch_add_explicit(&out.channel[2][rgb.b], 1, memory_order_relaxed);
atomic_fetch_add_explicit(&out.channel[3][Y], 1, memory_order_relaxed);

//
// Назад отдаем туже тектсуру что и прилетела.
//
outTexture.write(inColor,gid);
}

Для проектирования обвязки самого фильтра на Metal Framework воспользуемся уже знакомой оберткой DPCore3, скрывающая некоторые избыточные детали взаимоотношения с Metal.

Чтобы уж все совсем просто не было, создадим такую иерархию типов объектов для работы с гистограммой в общем виде:

  1. IMPHistogramBuffer — контейнер сырых данных для обмена между CPU и GPU. Будет просто структура с таблицей бинов.

  2. IMPHistogram — представление гистограммы в виде класса и операциями над данными гистограммы

  3. IMPHistogramAnalyzer — Metal-ический фильтр-анализатор данных связывающих GPU и CPU. Получает изображения или фреймы из камеры, выполняет контекст обработки и возвращает представление картинки в виде объекта IMPHistogram

  4. IMPHistogramSolver — протокол описания объектов решения конкретных задачи анализа, к таким солверам могут быть отнесены, к примеру, штука ищущая цветовую доминанту изображения, динамический диапазон картинки, диапазон интенсивностей и т.п…

Конструктор анализатора

В конструкторе анализатора связываем фильтр с нашей kernel-функцией счета и создаем в общей памяти CPU/GPU область которая будет хранить наш контейнер.

///
/// Конструктор анализатора с произвольным счетчиком, который
/// задаем kernel-функцией. Главное условие совместимость с типом IMPHistogramBuffer
/// как контейнером данных гистограммы.
///
///
init(function: String, context aContext: DPContext!) {
super.init(context: aContext)

// инициализируем счетчик
kernel_impHistogramCounter = DPFunction.newFunction(function, context: self.context)

// создаем память в устройстве под контейнер счета
histogramUniformBuffer = self.context.device.newBufferWithLength(sizeof(IMPHistogramBuffer), options: MTLResourceOptions.CPUCacheModeDefaultCache)

// добавляем счетчик как метод фильтра
self.addFunction(kernel_impHistogramCounter);
}

Обмен данными гистограммы с GPU

override func configureFunction(function: DPFunction!, uniform commandEncoder: MTLComputeCommandEncoder!) {
commandEncoder.setBuffer(histogramUniformBuffer, offset:0, atIndex:0);
}

И вот тут нам таки пригодились blit-операции над памятью GPU

Как можно заметить. в методе обмена данными между CPU/GPU по идее бы нужно было бы сделать, что-то вроде:

... memset(histogramUniformBuffer.contents(),0,sizeof(IMPHistogramBuffer))
...

Перед каждой установкой буфера данных. Логично же обнулить память до запуска нового расчета новой гистограммы. И вот тут всплывает одно интересное свойство Metal Framework, да мы можем так сделать, но запись значений в ячейки будет производится процессором, а это медленно, и не гарантирует корректной записи в цикле потока команд. И для таких вот случаев Apple-ом нам даден инструмент в виде работы с MTLBlitCommandEncoder позволяющий быстро и безконфликтно на аппаратном уровне выполнять операции над кусками памяти GPU и разделяемой памяти CPU/GPU, т.е. над буферами и текстурами.

///
/// При каждом обращении к GPU для расчета гистограмы нам нужно обресетить данные посчитанные на предыдущем этапе
/// если объект анализатора постоянно определен.
///
override func configureBlitUniform(commandEncoder: MTLBlitCommandEncoder!) {
commandEncoder.fillBuffer(histogramUniformBuffer, range: NSMakeRange(0, sizeof(IMPHistogramBuffer)), value: 0)
}

 

Что-то вроде заключения

В этом посте мы рассмотрели пример того, как легко, используя Metal, можно посчитать произвольную гистограмму произвольного изображения на GPU (Apple, где же ты была раньше!?) и дальше как-то её использовать. О том как использовать будет отдельный пост — в нем мы рассмотрим простые примеры вычисления на базе гистограммы доминантного цвета изображения для коррекции баланса белого и растяжения гистограммы изображения для повышения его контраста.   А пока полный исходный код новой версии Metalagram-а использующий вот это все, можно посмотреть: тут, в репозитории примеров проекта: ImageMetalling.


Еще раз напомню: авторы блога не преследуют задачи быть предельно корректным, но если заметили явную ашипку, если написали явную глупость, если что-то не понятно: комментируйте или пишите на: imagemetalling [*] gmail.com .

Реклама

Расчет гистограммы изображения на GPU. Атомарные типы данных Metal и операции над ними: 2 комментария

Добавить комментарий

Заполните поля или щелкните по значку, чтобы оставить свой комментарий:

Логотип WordPress.com

Для комментария используется ваша учётная запись WordPress.com. Выход / Изменить )

Фотография Twitter

Для комментария используется ваша учётная запись Twitter. Выход / Изменить )

Фотография Facebook

Для комментария используется ваша учётная запись Facebook. Выход / Изменить )

Google+ photo

Для комментария используется ваша учётная запись Google+. Выход / Изменить )

Connecting to %s