Доступ к пикселям текстуры в Metal

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

В последнее время я наблюдаю некоторый рост интереса к Metal со стороны разработчиков: переодически получаю отзывы и благодарности от вас, и даже ссылочки на приложения в AppStore, при разработке которых помогли посты этого ресурса. Это приятно, чоуж.

С другой стороны, Apple таки приближается к Swift 5, замораживает синтаксис и объявляет совместимость по ABI. И хотя Objc или C++ реализации приложений на Metal все еще субъективно шустрее чем Swift, а код не нужно переписывать от версии к версии практически с нуля, Swift начинает выглядеть как штука, на которой уже можно писать примеры для важных постов без вопросов к версии  XCode для сборки.

А еще мне хочется проверить работу шорткода [Gist] вместо кривого и дурацкого  [Code]. Ну вы поняли.

Сегодня не будет математики

До сих пор с помощью Метал мы решали практические задачи требующие больших вычислений. Как правило из предметной и очень прикладной области. Однако есть бооольшой класс проблем (в хорошем смысле этого слова) сугубо технического характера: без этих ваших формул. Вот, к примеру, совсем недавно поступил вопрос: а как прочитать цвет конкретных пикселей текстуры? Т.е. казалось бы простая и очевидная вещь, но вызывает трудность. Поэтому я пошарился по докам Apple, стековерфло и оказалось, что действительно вопрос не раскрыт — ибо для разработчика SDK это самоочевидно, но новичку не доступно.

Давайте сконструируем простое приложение

Приложение будет давать возможность показать в окошке картинку из файла. Мы сможем тыкнуть мышкой в картинку, нарисовать квадратную область, и определить средний цвет пикселов этой квадратной области. Размер области может быть и 1 пиксел, например. Но мы в коде зададим квадрат 20×20.

Традиционно исходники приложения в папочке ImageMetalling-14 если текст поста читать лень. Сборка как всегда с cocoapods:

$ pod install
$ open ImageMetalling-14.xcworkspace

И дальше в XCode запускаем проект.

Что под капотом?

Давайте начнем с того, что хостовый слой Metal предоставляет нам несколько базовых структур данных для взаимодействия с ядрами GPU: MTLTexture и MTLBuffer. По сути это обертки кусков памяти, которые мы можем размещать в областях CPU и GPU одновременно или копировать между различными типами памяти в контексте pipeline обработки. В зависимости от типа устройства при размещении данных в этих кусках нам нужно будет синхронизовать контекст с текущим тредом исполнения в CPU или нет. Но если опустить детали, то обе структуры — это просто последовательность данных, которую мы можем читать и изменять как в хостовом треде исполняемом на CPU так и на конкретном ядре GPU. Если мы что-то записали в текстуру, передали в ядро и произвели вычисления, то мы можем результаты вычислений разместить в другой кусок памяти, на который может быть линейно отображена произвольная структура. Вычисления могут быть записаны в результирующую текстуру, а может и в специальный объект.

По сути, MTLTexture отличается от MTLBuffer только тем, что текстура уже имеет специфический набор методов по доступу к её специально-структурированным данным зависящим от типа представления пикселей и семплов, а MTLBuffer — может содержать пользовательский тип данных с нужным разработчику набором полей и методов (ну и + какой-то заданный набор уже есть в обоих представлениях CPU и GPU).

Ну, мы так и поступим возьмем центр области, пробежимся по окружающим пикселям, сложим да поделим. А результат разместим в буфер с данными типа float3. Гениально же!

А, ну да, запустим столько ядер сколько нам нужно областей посчитать. Для задачи нам достаточно одного. Хотя это и из пушки по воробьям, но нам ехать, а не о культуре оптимизации.

Реализация ядра

//
// Ядро чтения семплов текстуры с точностью до индекса и вычисление среднего значения цвета областей с центрами
// в пространстве RGB
//
kernel void get_patchColors(
// исходная текстура
metal::texture2d<float, metal::access::sample> source [[texture(0)]],
// список центров в которых нужно проситать семплы
device float2 *centers [[ buffer(0) ]],
// список усредненных цветов областей в которых мы прочитаем тектсуру
device float3 *colors [[ buffer(1) ]],
// размер квадратной области
constant float &regionSize [[ buffer(2) ]],
// позиция треда GPU в гриде == индекс центра области
uint2 tid [[thread_position_in_grid]]
)
{
uint width = source.get_width();
uint height = source.get_height();
float2 size = float2(width,height);
float2 point = centers[tid.x];
int rs = -regionSize/2;
int re = regionSize/2+1;
uint2 gid = uint2(float2(point.x,point.y) * size);
colors[tid.x] = avrgColor(rs, re, rs, re, gid, source);
}

Усреднение цвета

///
/// Расчитываем средний цвет области текстуры
///
inline float3 avrgColor(int startx, // начало области по x
int endx, // конец области по x
int starty, // начало по y
int endy, // конец y
uint2 gid, // индекс семпла
texture2d<float> source // текстура
){
float3 color(0);
float3 c(0);
for(int i = startx; i<endx; i++ ){
for(int j = starty; j<endy; j++ ){
uint2 gid2 = uint2(int2(gid)+int2(i,j));
float3 s = source.read(gid2).rgb;
color += s;
c+=float3(1);
}
}
return color/c;
}
view raw avrgColor.metal hosted with ❤ by GitHub

Boilerplate и вот это все в Metal

Тут должен быть TL;DR про то, что на чистом Metal кодить хоть и сильно удобнее по сравнению с OpenGL и даже OpenCL, да и Cuda иногда дает прикурить. Но для задач писания процессинга изображений все еще нужно кричать много слов и складывать кучу букв в осмысленный код. Поэтому для себя у нас есть обёрточка уменьшающая количество страдания и клавиатурной боли. Как полагается, в общем.

  • SDK IMProcessing — обертка для классов работы с фильтрами и их конструированием для нужд процессинга и анализа изображений;
  • SDK IMProcessingUI — обертка для объектов рендеринга на экран результатов фильтрации и анализа изображений.

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

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

/// Установка областей чтения семплов
public var centers:[float2] = [float2]() {
didSet{
if centers.count > 0 {
// создать MTL-буфер из которого будем читать центры областей в шейдере
centersBuffer = context.device.makeBuffer(length: MemoryLayout<float2>.size * centers.count, options: [])!
// создаем MTL-буфер в который будем писать значения цветов в шейдере
colorsBuffer = context.device.makeBuffer(length: MemoryLayout<float3>.size * centers.count, options: .storageModeShared)!
// пишем в буфер центры
memcpy(centersBuffer.contents(), centers, centersBuffer.length)
// выделяем память массив цветов, в который потом скопируем то что прилетело
// в буфер в шейдере
_colors = [float3](repeating:float3(0), count:centers.count)
// определим размерность грида вычислений GPU
patchColorsKernel.preferedDimension = MTLSize(width: centers.count, height: 1, depth: 1)
// сбросим фильтр
dirty = true
// запустим вычисления на шейдерах
process()
}
}
}

Передаем параметры в кернел:

if self.centers.count > 0 {
// центры - читаем
command.setBuffer(self.centersBuffer, offset: 0, index: 0)
// цвета - пишем
command.setBuffer(self.colorsBuffer, offset: 0, index: 1)
// размер области
command.setBytes(&self.regionSize, length:MemoryLayout.size(ofValue: self.regionSize), index:2)
}

Читаем буфер с цветами в массив:

if self.centers.count > 0 {
// копирем в цвета данные из буфера GPU в массив цветов
memcpy(&self._colors, self.colorsBuffer.contents(), self.colorsBuffer.length)
}

Полный кода класса:  ColorObserver.swift. Он же в исходном SDK IMPColorObserver.

Отрисовка цвета в окошке приложения

/// Создаем фильтр обзёрвера цветов текстуры
private lazy var patchColors:ColorObserver = {
let f = ColorObserver(context: self.context)
//
// Размер прямоугольной (квадратной) области по которой мы интерполируем
// (на самом деле усредняем) цвет текстуры
//
f.regionSize = 20
//
// Добавляем к фильтру обработку событий пересчета целевой тектстуры,
// которая на самом деле не пересчитывается и читает в шейдере в буфер её RGB-смеплы
//
f.addObserver(destinationUpdated: { (destination) in
//
// Поскольку мы читаем только одну область то берем первый элемент массива
// прочитаных семполов цветов
//
var rgb = f.colors[0]
// представление [0-1] в NSColor
let color = NSColor(color: float4(rgb.r,rgb.g,rgb.b,1))
// инвертируем цвет
let inverted_rgb = float3(1) - rgb
let inverted_color = NSColor(color: float4(inverted_rgb.r,inverted_rgb.g,inverted_rgb.b,1))
// для отображения в textfield переведем в 8-битное представление
rgb = rgb * float3(255)
DispatchQueue.main.async {
// просто рисуем
self.patch.strokeColor = inverted_color
self.colorLabel.backgroundColor = color
self.colorLabel.stringValue = String(format: "%3.0f, %3.0f, %3.0f", rgb.r, rgb.g, rgb.b)
}
})
return f
}()

Тут мы просто ловим факт асинхронного завершения исполнения процессинга на GPU и отрисовываем результат уже традиционными средствами OSX.

 

Выводы

Бок есть! И он либо правый либо левый. Но иногда надо посмотреть снизу.

 


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

Доступ к пикселям текстуры в Metal: Один комментарий

Оставьте комментарий