OpenCL Kernel работает медленнее на более быстром GPU - PullRequest
Новая
       43

OpenCL Kernel работает медленнее на более быстром GPU

0 голосов
/ 25 января 2012

Я довольно новичок в OpenCL и пытался реализовать алгоритм DSP, чтобы сравнить его производительность на разных графических процессорах по сравнению со стандартной реализацией процессора.Хотя я добился огромного прироста производительности, я нахожу странным то, что на GT240 я получаю почти такой же прирост, как и гораздо более быстрый GTX 480. Моя программа выполняет два ядра, и пока одно ускоряется на GTX 480, другое замедляетсявниз.

GT240: Ядро 1: 226us, Ядро 2: 103us.

GTX 480: Ядро 1: 35us, Ядро 2: 293us.

Эти числа были получены с использованиемВизуальный Профилировщик.Ниже приведен код для Kernel 2, который почти в 3 раза медленнее на большой карте.Это ядро ​​берет блок памяти размером iTotalBins x iNumAngles и вычисляет максимум для каждой строки длины iNumAngles, а также подгоняет кривую к трем смежным значениям. 

__kernel void max_curve_fit_gpu (__global float* fCorrelationResult,
                          const int iNumAngles,
                          const int iTotalBins,
                          __global float* fResult){

// Get the thread ID which is used as the index the bin the direction is being calculated for
const int iBinNum = get_global_id(0);
const int iOffset = iBinNum*iNumAngles;

// Find the max for this bin
float fMax = 0;
int iMaxIndex = 0;
for (int iAngle=0; iAngle<iNumAngles; iAngle++)
{
    if (fMax < fCorrelationResult[iOffset + iAngle])
    {
        fMax = fCorrelationResult[iOffset + iAngle];
        iMaxIndex = iAngle;
    }
}

// Do the curve fit
float fPrev, fNext, fA, fB, fAxis;
fPrev = fCorrelationResult[iOffset + (iMaxIndex + iNumAngles - 1) % iNumAngles];
fNext = fCorrelationResult[iOffset + (iMaxIndex + iNumAngles + 1) % iNumAngles];

fB = (fPrev - fNext)*0.5f;
fA = (fNext + fPrev) - fMax*2.0f;
fAxis = fB / fA;

    // Store the result
fResult[iBinNum] = iMaxIndex + fAxis; }

Visual Profiler также указывает, чтосоставляет 135% Воспроизведение инструкций глобальной памяти для ядра 2. У меня есть версия max search, в которой не используется оператор if-else, но он работает еще медленнее на обоих графических процессорах.

Любая помощь будет принята с благодарностью.

Ответы [ 2 ]

3 голосов
/ 26 января 2012

В вашем коде поток T получит доступ к fCorrelationResult[T*iNumAngles+iAngle], что означает, что у вас нет объединенных обращений и, возможно, конфликты банка памяти тоже Банковские конфликты могут объяснить явление, которое вы наблюдаете.

Вы должны транспонировать свою матрицу и получить доступ к fCorrelationResult[T+iAngle*iNumBins] вместо этого. Вы, безусловно, получите хорошее ускорение и, возможно, более регулярные тесты между двумя графическими процессорами.

0 голосов
/ 25 января 2012

С OpenCL, возможно ли перейти на более низкий уровень и узнать что-нибудь об использовании регистров и общей памяти при работе ядра на определенном графическом процессоре?

Из-за моего ограниченного знакомства с NVIDIA CUDA, использование может быть ключевым здесь. GT240 имеет вычислительные возможности 1.2, а GTX480 - 2.0, поэтому последний имеет 2 регистра и 3x разделяемой памяти. Я предполагаю, что код, созданный OpenCL для второго ядра, не может использовать эти ресурсы на 480-м. Может возникнуть конфликт банка с общей памятью, например.

...