Производительность записи в память - общая память процессора GPU

Я выделяю как входные, так и выходные данные MTLBuffer с помощью posix_memalign в соответствии с совместной графической/процессорной документацией, предоставленной memkite.

Кроме того, проще всего использовать новейший API, чем гадать с posix_memalign

let metalBuffer = self.metalDevice.newBufferWithLength(byteCount, options: .StorageModeShared)

Моя функция ядра работает примерно с 16 миллионами сложных структур значений и записывает в память равное количество структурных структур с целыми значениями.

Я выполнил некоторые эксперименты, и моя сложная математическая секция "Ядро метала" выполняется за 0.003 секунды (Да!), но запись результата в буфер занимает > 0.05 (Нет!) секунд. В моем эксперименте я прокомментировал математическую часть и просто назначил нуль памяти, и она занимает 0.05 секунды, комментируя назначение и добавляя обратно математику, 0.003 секунды.

В этом случае медленная разделяемая память, или есть какой-то другой совет или трюк, который я могу попробовать?

Дополнительная информация

Тестовые платформы

iPhone 6S - ~ 0.039 секунд за кадр
iPad Air 2 - ~ 0.130 секунд за кадр

Данные потоковой передачи

Каждое обновление шейдера получает приблизительно 50000 комплексных чисел в виде пары типов float в структуре.

struct ComplexNumber {
    float real;
    float imaginary;
};

Подпись ядра

kernel void processChannelData(const device Parameters *parameters [[ buffer(0) ]],
                               const device ComplexNumber *inputSampleData [[ buffer(1) ]],
                               const device ComplexNumber *partAs [[ buffer(2) ]],
                               const device float *partBs [[ buffer(3) ]],
                               const device int *lookups [[ buffer(4) ]],
                               device float *outputImageData [[ buffer(5) ]],
                               uint threadIdentifier [[ thread_position_in_grid ]]);

Все буферы содержат - в настоящее время - неизменные данные, кроме inputSampleData, которые получают 50000 выборок, над которыми я буду работать. Другие буферы содержат примерно 16 миллионов значений (128 каналов х 130000 пикселей) каждый. Я выполняю некоторые операции над каждым "пикселем" и суммирую сложный результат по каналам и, наконец, беру абсолютное значение комплексного числа и назначая полученный результат float на outputImageData.

Отправка

commandEncoder.setComputePipelineState(pipelineState)

commandEncoder.setBuffer(parametersMetalBuffer, offset: 0, atIndex: 0)
commandEncoder.setBuffer(inputSampleDataMetalBuffer, offset: 0, atIndex: 1)
commandEncoder.setBuffer(partAsMetalBuffer, offset: 0, atIndex: 2)
commandEncoder.setBuffer(partBsMetalBuffer, offset: 0, atIndex: 3)
commandEncoder.setBuffer(lookupsMetalBuffer, offset: 0, atIndex: 4)
commandEncoder.setBuffer(outputImageDataMetalBuffer, offset: 0, atIndex: 5)

let threadExecutionWidth = pipelineState.threadExecutionWidth
let threadsPerThreadgroup = MTLSize(width: threadExecutionWidth, height: 1, depth: 1)
let threadGroups = MTLSize(width: self.numberOfPixels / threadsPerThreadgroup.width, height: 1, depth:1)

commandEncoder.dispatchThreadgroups(threadGroups, threadsPerThreadgroup: threadsPerThreadgroup)
commandEncoder.endEncoding()
metalCommandBuffer.commit()
metalCommandBuffer.waitUntilCompleted()

Пример GitHub

Я написал пример под названием Slow и разместил его на GitHub. Кажется, узким местом является запись значений во входной буфер. Итак, я думаю, вопрос заключается в том, как избежать узкого места?

Копия памяти

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

Текущее состояние

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

Ответ 1

Уменьшить размер типа

Первоначально я предварительно конвертировал подписанные 16-разрядные Целочисленные числа как Float (32-разрядные), так как в конечном итоге именно так они будут использоваться. Это случай, когда производительность начинает заставлять вас хранить значения в виде 16 бит, чтобы сократить размер данных пополам.

Objective-C над Swift

Для кода, связанного с перемещением данных, вы можете выбрать Objective-C по Swift (рекомендация Уоррен Мур). Производительность Swift в этих особых ситуациях все еще не до нуля. Вы также можете попробовать обратиться к memcpy или аналогичным методам. Я видел несколько примеров, которые использовали для буферов-указателей для цикла, и это в моих экспериментах выполнялось медленно.

Сложность тестирования

Я действительно хотел сделать некоторые из экспериментов по отношению к различным методам копирования на игровой площадке на машине, и, к сожалению, это было бесполезно. Версия iOS для одних и тех же экспериментов выполнялась совершенно по-разному. Можно подумать, что относительная производительность будет аналогичной, но я обнаружил, что это тоже недопустимое предположение. Было бы очень удобно, если бы у вас была игровая площадка, которая использовала устройство iOS в качестве интерпретатора.