Использование GPU для преобразования в пиксельные форматы?

Question

Я пытаюсь очень эффективно конвертировать упакованные изображения RGB 24bpp в упакованный RGBA 32bpp. Я пробовал использовать vImageConvert_RGB888toRGBA8888 от Accelerate.framework, но задавался вопросом, был ли более быстрый способ, используя вычислительное ядро ​​в Metal. Я пробовал несколько разных подходов в Metal, но результат всегда значительно медленнее, чем с Accelerate.framework, даже для больших изображений с> 1M пикселей.

Вот что мое ядро ​​вычисления выглядит следующим образом:

kernel void rgb24_to_rgba32(texture2d<half, access::read> inTexture [[texture(0)]], 
        texture2d<half, access::write> outTexture [[texture(1)]], 
        uint2 id [[ thread_position_in_grid ]]) 
{ 
    uint2 srcAddr1 = uint2(id.x * 3, id.y); 
    uint2 srcAddr2 = uint2(id.x * 3 + 1, id.y); 
    uint2 srcAddr3 = uint2(id.x * 3 + 2, id.y); 

    outTexture.write(half4(inTexture.read(srcAddr1).r, inTexture.read(srcAddr2).r, inTexture.read(srcAddr3).r, 1), id); 

    return; 
} 

Я определения inTexture как r8Unorm и outTexture как bgra8Unorm. Обе текстуры загружаются с использованием .storageModeShared, поэтому не должно быть никаких копий памяти.

Код работает, и преобразование выполняется правильно, но производительность не впечатляет. Я пробовал разные настройки threadgroupsPerGrid и threadsPerThreadgroup, но ни один из них не достиг сопоставимой производительности до Accelerate.framework.

Например, на A7 (iPad первого поколения) изображение 1024x1024 занимает около 32 мс, по сравнению с 6 мс, используя Accelerate.framework. Интересно, что разница намного меньше для более быстрого устройства, такого как iPhone 6 на базе A9 (1,5 мс на графическом процессоре против 1,1 мс с использованием Accelerate), но реализация Металла всегда медленнее.

Это просто не совместимая с GPU операция (возможно, из-за бесчисленных нерациональных доступов к памяти?) Могу ли я упустить что-то фундаментальное с точки зрения максимизации производительности моего вычислительного ядра?

UPDATE: Я был в конечном итоге удалось достичь значительно более высокую производительность, чем описанный выше, с использованием следующей реализации:

Этот подход использует 96-битный читает с помощью packed_uint3, и 128-битный записывает с помощью packed_uint4, чтобы значительно повысить производительность ,

#define RGB24_TO_RGBA32_PIXEL1(myUint) (myUint | 0xff000000) 

#define RGB24_TO_RGBA32_PIXEL2(myUint1, myUint2) (myUint1 >> 24 | \ 
               ((myUint2) << 8) | 0xff000000) 


#define RGB24_TO_RGBA32_PIXEL3(myUint2, myUint3) (myUint2 >> 16 | \ 
               ((myUint3) << 16) | 0xff000000) 

#define RGB24_TO_RGBA32_PIXEL4(myUint3) ((myUint3 >> 8) | 0xff000000) 

inline packed_uint4 packed_rgb24_to_packed_rgba32(packed_uint3 src) { 
    return uint4(RGB24_TO_RGBA32_PIXEL1(src[0]), 
       RGB24_TO_RGBA32_PIXEL2(src[0], src[1]), 
       RGB24_TO_RGBA32_PIXEL3(src[1], src[2]), 
       RGB24_TO_RGBA32_PIXEL4(src[2])); 
} 

kernel void rgb24_to_rgba32_textures(
         constant packed_uint3 *src [[ buffer(0) ]], 
         device packed_uint4 *dest [[ buffer(1) ]], 
         uint2 id [[ thread_position_in_grid ]]) 
{ 
    // Process 8 pixels per thread (two packed_uint3s, each containing 4 pixels): 
    uint index = id.x * 2; 
    dest[index] = packed_rgb24_to_packed_rgba32(src[index]); 
    dest[index + 1] = packed_rgb24_to_packed_rgba32(src[index + 1]); 
    return; 
} 

При таком подходе, дифференциальная производительность на старых устройствах становится гораздо меньше (ускорение примерно 2 раза быстрее, чем GPU), а на более современных устройств (A9), металла на самом деле ветры быть около 40-50 % быстрее.

Я пробовал обрабатывать один, два или более packed_uint3 векторов в потоке, и вывод заключался в том, что два вектора - это сладкое пятно для производительности.

Answer 1

Только ради закрытия, вот ответ от разработчиков Apple на этот вопрос. Суть в том, что GPU просто не предлагает никаких реальных преимуществ в этом случае, потому что это преобразование не является тяжелой операцией с вычислением.

После обсуждения с инженерной и оценки больше образцов реализации, приговор на Metal В.С. Ускорьте производительность для преобразования упакованных изображений RGB 24bpp в упакованные 32bpp Изображения RGBA: на новых устройствах вы можете приблизиться к тем же , используя Metal, но ускорение будет быстрее для этой операции . «VImage чрезвычайно хорошо настроенное внедрение и так эта операция преобразования не вычислит тяжелой лучшее, что мы можем сделать, это быть в паритете.»

Предложенное обоснованием этого локальность данных и эффективен работает на несколько пикселов за раз (что вы упомянули). Самая быстрая протестированная металлическая реализация обработала два пикселя на нить и все еще отставала от vImageConvert_RGB888toRGBA8888.

Было реализовано «оптимизированное» исполнение с использованием металлических буферов, а не , чем текстуры (что-то еще, что вы упомянули об изучении) и . На удивление этот подход был несколько менее впечатляющим.

Наконец, в обсуждение вошли настройки групп нитей, а также настройка путем добавления кода в ядро ​​для обработки случая, когда положение нити в сетке находится за пределами целевого изображения. Опять же, , несмотря на эти соображения Ускорение оставалось как быстрая реализация .

Я должен добавить, что одним из реальных преимуществ использования металла является использование ЦП, в то время как оно не ускоряется, это значительно снижает рабочую нагрузку процессора. Для приложений, где процессор сильно загружен, подход Metal может иметь смысл.

Answer 2

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

• Рассмотрите возможность использования буферов вместо текстур. Ваш буфер ввода может быть типа packed_char3, и ваш выходной буфер может быть типа packed_char4. Затем, вместо того, чтобы делать три чтения текстуры для каждой записи, вы можете индексировать в исходный буфер только один раз на пиксель. Как вы заметили, большинство этих чтений будут неравнозначными, но этот подход может сэкономить вам некоторые конверсии формата и пропускную способность.

• Рассмотрите возможность выполнения дополнительной работы при вызове ядра.Если ваши размеры изображения кратно 4 или 8 (например), вы можете использовать цикл (который должен быть разворачиван компилятором) для обработки этого количества пикселей в ядре, тем самым уменьшая количество групп потоков, которые вам нужно отправить.

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