CUDA not so fast against CPU with OpenMP?

Question

I am trying to compute cross-correlation amongst 450 vectors each of size 20000. While doing this on CPU i stored the data in 2D matrix with rows=20000 and cols=450.

The serial code for the computation looks like

    void computeFF_cpu( float * nSamples, float * nFeatures, float ** data, float ** corr
        #pragma omp parallel for shared(corr, data)
        for( int i=0 ; i<nFeatures ; i++ )
        {
            for( int j=0 ; j<nFeatures ; j++ )
                corr[i][j] = pearsonCorr( data[i], data[j], nSamples );
        }

int main()
{
.
.
**for( int z=0 ; z<1000 ; z++ )**
computeFF_cpu( 20000, 450, data, corr );
.
.
}

This works perfectly. Now I have attempted to solve this problem with GPU. I have converted the 2D data matrix into row-major format in GPU memory and I have verified that the copy is correctly made.

The vectors are stored as a matrix of size 900000 (ie. 450*20000) in row major format. Organized as follows
<---nSamples of f1---><---nSamples of f2 ---><---nSamples of f3--->......

My cuda code to compute cross-correlation is as follows

    // kernel for computation of ff
    __global__ void computeFFCorr(int nSamples, int nFeatures, float * dev_data, float * dev_ff)
    {
        int tid = blockIdx.x + blockIdx.y*gridDim.x;
        if( blockIdx.x == blockIdx.y )
        dev_ff[tid] = 1.0;
        else if( tid < nFeatures*nFeatures )
        dev_ff[tid] = pearsonCorrelationScore_gpu( dev_data+(blockIdx.x*nSamples), dev_data+(blockIdx.y*nSamples), nSamples );
    }

    main()
    {
    .
    .
        // Call kernel for computation of ff
**for( int z=0 ; z<1000 ; z++ )**
        computeFFCorr<<<dim3(nFeatures,nFeatures),1>>>(nSamples, nFeatures, dev_data, corr);
        //nSamples = 20000
        // nFeatures = 450
        // dev_data -> data matrix in row major form
        // corr -> result matrix also stored in row major
    .
    .
    }

Are you sure that "the inefficiency that may be caused by having just 1 thread per block" is really ignorable in this case? Does the code actually run to completion if you use a smaller input space? — talonmies
well, program completes in about 2min. How do you suggest to parallize into N threads (say) in sample? I am not able to think really. I am a newbie to cuda — mkuse
well, that clearly answers your question. If there was deadlock, the code would never finish. There is no deadlock. It is just very slow. I would start by fixing the question to reflect the true nature of the problem. — talonmies
@talonmies What more information do you want in this regard? What I am trying to accomplish here is feature-feature correlation. What I mean is that, I have a 2D matrix of size 20000x450. Where each row represents a vector. I need to find another matrix which gives correlation of each pair of vectors. — mkuse
I mean your entire question is premised on the idea that your code never finishes or has some kind of deadlock. But that clearly isn't the case. You questions is "what's wrong with this, is it deadlocking?". And the answer is "nothing wrong, just slow, and you already know why". If you are really meaning to ask a different question, then by all means go ahead and edit your question. — talonmies

mkuse mkuse · Accepted Answer · 2013-02-01T13:07:00

Seems like I have found an answer to my own question. I have the following experiment. I have changed the values of z (ie. number of times the function is executed). This kind of approach was suggested in several of previous post on stackoverflow under cuda tag.

Here is the table --

Z=100 ; CPU=11s ; GPU=14s
Z=200 ; CPU=18s ; GPU=23s
Z=300 ; CPU=26s ; GPU=34s
Z=500 ; CPU=41s ; GPU=53s
Z=1000; CPU=99s ; GPU=101s
Z=1500; CPU=279s; GPU=150s
Z=2000; CPU=401s; GPU=203s

It is evident that as the number of computation grows GPU is able to scale much better than the CPU.

CUDA not so fast against CPU with OpenMP?

1 Answers