OpenCL Kernel performs slower on faster GPU

Question

I am fairly new to OpenCL and have been attempting to implement a DSP algorithm to compare its performance on different GPU's compared to the standard CPU implementation. Though I have achieved a massive performance gain, what I find strange is that I get the almost the same gain on GT240 as a much faster GTX 480. My program executes two kernels, and while the one speeds up on the GTX 480 the other slows down.

GT240: Kernel 1: 226us, Kernel 2: 103us.

GTX 480: Kernel 1: 35us, Kernel 2: 293us.

These numbers were obtained using Visual Profiler. Below is the code for Kernel 2, which is almost 3 times slower on the bigger card. This kernel takes a block of memory which is iTotalBins x iNumAngles big, and computes the max for each row of iNumAngles length, and fits a curve to the 3 adjacent values.

__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 also indicates that there is 135% Global memory instruction replay for Kernel 2. I have a version of the max search which does not use a if-else statment, but it runs even slower on both GPU's.

Any help will be greatly appreciated.

Eric Bainville Eric Bainville · Accepted Answer · 2012-01-25T22:59:08

In your code, thread T will access fCorrelationResult[T*iNumAngles+iAngle], meaning you have no coalesced accesses, and probably memory bank conflicts too. Bank conflicts may explain the phenomenon you observe.

You should transpose your matrix, and access fCorrelationResult[T+iAngle*iNumBins] instead. You will certainly get a nice speedup, and probably more regular benchmarks between the two GPU's.

OpenCL Kernel performs slower on faster GPU

2 Answers