Повышение качества изображения шейдеров Metal Performance для iOS

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

2 5 6
3 6 7
8 9 0

изображение с повышенным разрешением должно быть

2 2 5 5 6 6
2 2 5 5 6 6
3 3 6 6 7 7
3 3 6 6 7 7
8 8 9 9 0 0
8 8 9 9 0 0

Я написал следующую функцию ядра:

kernel void upsample(texture2d<float, access::read> inTexture [[texture(0)]],
                     texture2d<float, access::write> outTexture [[texture(1)]],
                     uint2 gid [[thread_position_in_grid]])
{
    const float4 color = inTexture.read(gid);

    uint2 p;
    p = uint2(gid.x * 2, gid.y * 2);
    outtexture.write(color, p);
    p = uint2(gid.x * 2 + 1, gid.y * 2); 
    outtexture.write(color, p);
    p = uint2(gid.x * 2, gid.y * 2 + 1);
    outtexture.write(color, p);
    p = uint2(gid.x * 2 + 1, gid.y * 2 + 1);  
    outtexture.write(color, p);
}

Но я не уверен, правильно ли это. Как мне узнать, что исходный "gid" связан с координатами входной текстуры, но не с выходной текстурой?


person fireman    schedule 17.09.2016    source источник
comment
В качестве примечания: если вы хотите максимизировать производительность с большими наборами данных, мой опыт подсказывает, что вы можете значительно улучшить производительность, используя для этого фиксированное оборудование графического процессора. Просто настройте конвейер рендеринга, загрузите текстуру, настройте целевой рендеринг должного размера, установите сэмплер текстуры на «ближайшее» увеличение и позвольте фиксированному оборудованию выполнять масштабирование. В других подобных тестах, которые я проводил, я видел, как фиксированное оборудование делает подобные вещи в несколько раз быстрее, чем вычислительное ядро.   -  person ldoogy    schedule 30.10.2016


Ответы (3)


То, что вы написали, мне кажется правильным, если предположить, что вы отправили сетку, имеющую те же размеры, что и исходная текстура.

Нет никакого внутреннего соответствия между положением потока в сетке и координатами, которые вы читаете или записываете. gid - это просто точка в таблице рабочих элементов, которые вы отправили с помощью кодировщика вычислительных команд. В вашем ядре вы можете интерпретировать это как хотите. Если вы отправляете 2D-сетку и работаете с текстурами, часто имеет смысл рассматривать gid как пару координат, как вы это делаете здесь.

Вы также могли бы отправить одномерную сетку и «обернуть» ее в двухмерную сетку самостоятельно, но использование позиции двухмерного потока может сэкономить вам (дорогостоящее) целочисленное деление в ядре, помимо того, что оно будет более интуитивно понятным.

person warrenm    schedule 17.09.2016

@fireman Я профилирую ваш код выборки, но запись 4 данных один за другим обходится дорого. И твой вопрос про гид. это связано с рабочими элементами вашей области процесса ядра для текстуры ввода и вывода. Может быть, он содержит 1 пиксель на входе и содержит 4 пикселя для вывода.

person Ericking    schedule 09.01.2017

Теперь эта потребность решается семейством ядер MPSCNNUpsampling. См. MetalPerformanceShaders / MPSCNNUpsampling.h

person Ian Ollmann    schedule 15.08.2017
comment
Я знаю, что это было в августе, но это не обязательно делает работу, как будто OP хочет использовать только MTLTexture, это может быть дорогостоящей операцией, чтобы выполнить преобразование только для повышения дискретизации. - person DaveNine; 19.12.2017