Непонятны результаты профилирования моего ядра OpenCL (окнная функция DSP)
-
30-09-2019 - |
Вопрос
Я закончил ядро оконной функции в OpenCL.По сути, оконная функция просто применяет набор коэффициентов к другому набору чисел по частям (Википедия объясняет это лучше).В большинстве случаев мне удавалось поместить массив с плавающей точкой оконных коэффициентов в кеш констант.
Я ожидал, что результаты Compute Prof покажут, что передача памяти между хостом и устройством между хостами займет более 95% времени обработки.Почти во всех моих случаях это составляет только 80% времени обработки.Я записываю и считываю один массив с плавающей запятой на 4,2 миллиона на плату и обратно, а также пишу другой массив с плавающей запятой, который обычно остается значительно ниже миллиона.
Что-нибудь в ядре выглядит подозрительно?Есть какие-либо мнения о том, является ли эта проблема проблемой, которая должна работать быстрее на графическом процессоре, чем на процессоре (я все еще не на 100% уверен в этом).Я немного ошеломлен тем, почему мои gld_efficiency и gst_efficiency колеблются между 0,1 и 0,2.Я создал это ядро с учетом объединения глобальной памяти G80.Общая пропускная способность моей глобальной памяти кажется нормальной и составляет 40 ГБ.Ядро довольно простое и опубликовано ниже.
__kernel void window(__global float* inputArray, // first frame to ingest starts at 0. Sized to nFramesToIngest*framesize samples
__constant float* windowArray, // may already be partly filled
int windowSize, // size of window frame, in floats
int primitivesPerDataFrame, //amount of primitives in each frame of inputArray parameter
int nInFramesThisCall, //each thread solves a frame, so this integer represent how many threads this kernel launches
int isRealNumbers //0 for complex, non-zero for real
)
{
int gid = get_global_id(0) + get_global_size(0) * get_global_id(1);
if(gid < nInFramesThisCall) //make sure we don't execute unnecessary threads
{
if(isRealNumbers)
{
for(int i = 0; i < primitivesPerDataFrame; i++)
{
int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize];
}
}
else //complex
{
for(int i = 0; i < primitivesPerDataFrame; i++)
{
int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize/2];
}
}
}
}
Решение
Сколько потоков (кстати, термин OpenCL — рабочие элементы) вы используете?Вам нужно как минимум что-то из сотен, чтобы эффективно загрузить большой графический процессор.
Вы говорите, что хотите использовать объединенный доступ к памяти, но загрузка со смещением типа
int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
не сделает это возможным в большинстве случаев.NVidia G80 имеет довольно серьезные ограничения, когда дело доходит до объединения; дополнительную информацию см. в «Руководстве по передовому опыту OpenCL».По сути, рабочие элементы из одного варпа должны одновременно определенным образом обращаться к элементам выровненного блока размером 64 или 128 байт, чтобы обеспечить объединение операций загрузки и сохранения.
Или приведу вам пример:если primitivesPerDataFrame
равно 16, загрузка и сохранение варпа выполняются со смещениями, расположенными на расстоянии 16 элементов друг от друга, что делает любое эффективное объединение невозможным.