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

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

1 голос
/ 09 ноября 2010

Я закончил ядро ​​оконной функции в OpenCL.По существу, оконная функция просто применяет набор коэффициентов к другому набору чисел по частям (Википедия объясняет это лучше).В большинстве случаев мне удалось заполнить массив с плавающей запятой коэффициентом окна в постоянном кэше.

Я ожидал, что мои результаты от Compute Prof покажут, что передача данных с хоста на устройство и с хоста на хост займет более 95%время обработки.Почти во всех моих случаях это только 80% времени обработки.Я пишу и читаю один массив с плавающей запятой 4,2 миллиона на плате и с доски и пишу другой массив с плавающей точкой, который обычно остается значительно ниже миллиона.

Что-то в ядре выглядит подозрительно?Любые мнения о том, если это проблема, которая должна работать быстрее на GPU, чем на CPU в первую очередь (я все еще не на 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];
        }
    }
}

}

1 Ответ

1 голос
/ 11 ноября 2010

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

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

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

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

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

...