OpenCL slow - не знаю почему

Question

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

Я сосредоточен на вычислении векторов автокорреляции в OpenCL в качестве своей первой области улучшения. Код ЦП:

for (int i = 0; i < NrOfChannels; i++) { 
    for (int shift = 0; shift <= PredOrder[ChannelFilter[i]]; shift++) 
     vDSP_dotpr(Signal[i] + shift, 1, Signal[i], 1, &out, NrOfChannelBits - shift); 
} 
NrOfChannels = 6 
PredOrder = 129 
NrOfChannelBits = 150528. 

В моем тестовом файле эта функция занимает приблизительно 188 мс для завершения.

Вот мой метод OpenCL:

kernel void calculateAutocorrelation(size_t offset, 
           global const float *input, 
           global float *output, 
           size_t size) { 
size_t index = get_global_id(0); 
size_t end = size - index; 
float sum = 0.0; 

for (size_t i = 0; i < end; i++) 
    sum += input[i + offset] * input[i + offset + index]; 

output[index] = sum; 
} 

Вот как это называется:

Согласно приборам, моя реализация OpenCL, кажется, занимает около 13 мс, при этом около 54ms накладных расходов копирования памяти (gcl_memcpy).

Когда я использую гораздо больший тестовый файл, 1 минута 2-канальной музыки против 1 секунды 6-канального, в то время как измеренная производительность кода OpenCL кажется одинаковой, загрузка процессора падает примерно до 50 %, и вся программа занимает примерно 2 раза дольше.

Я не могу найти причины для этого в Инструментах, и я еще ничего не читал, что говорит о том, что я должен ожидать очень тяжелого перебоя в работе и выходе из OpenCL.

Answer 1

Если я правильно читаю код ядра, каждый рабочий элемент выполняет итерирование по всем данным из его местоположения до конца. Это не будет эффективным. Для одной (и основной проблемы производительности) доступ к памяти не будет объединен и поэтому не будет иметь полную пропускную способность памяти. Во-вторых, поскольку каждый рабочий элемент имеет разный объем работы, в рабочей группе будет разброс ветвей, что приведет к тому, что некоторые потоки ожидают других.

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

Чтобы узнать, как считывается память, выясните, как 16 рабочих элементов (например, global_id от 0 до 15) будут считывать данные для каждого шага.

Обратите внимание, что если каждый рабочий элемент в рабочей группе получает доступ к одной и той же памяти, существует оптимизация «трансляции», которую может произвести оборудование. Таким образом, просто изменение порядка вашего цикла может улучшить ситуацию.