Непонятны результаты профилирования моего ядра OpenCL (окнная функция DSP)

Question

Я закончил ядро ​​оконной функции в 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];
        }
    }
}

}

Answer 1

Сколько потоков (кстати, термин OpenCL — рабочие элементы) вы используете?Вам нужно как минимум что-то из сотен, чтобы эффективно загрузить большой графический процессор.

Вы говорите, что хотите использовать объединенный доступ к памяти, но загрузка со смещением типа

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

не сделает это возможным в большинстве случаев.NVidia G80 имеет довольно серьезные ограничения, когда дело доходит до объединения; дополнительную информацию см. в «Руководстве по передовому опыту OpenCL».По сути, рабочие элементы из одного варпа должны одновременно определенным образом обращаться к элементам выровненного блока размером 64 или 128 байт, чтобы обеспечить объединение операций загрузки и сохранения.

Или приведу вам пример:если primitivesPerDataFrame равно 16, загрузка и сохранение варпа выполняются со смещениями, расположенными на расстоянии 16 элементов друг от друга, что делает любое эффективное объединение невозможным.