6

I've written a CUDA4 Bayer demosaicing routine, but it's slower than single threaded CPU code, running on a16core GTS250.
Blocksize is (16,16) and the image dims are a multiple of 16 - but changing this doesn't improve it.

Am I doing anything obviously stupid?

--------------- calling routine ------------------
uchar4 *d_output;
size_t num_bytes; 

cudaGraphicsMapResources(1, &cuda_pbo_resource, 0);    
cudaGraphicsResourceGetMappedPointer((void **)&d_output, &num_bytes, cuda_pbo_resource);

// Do the conversion, leave the result in the PBO fordisplay
kernel_wrapper( imageWidth, imageHeight, blockSize, gridSize, d_output );

cudaGraphicsUnmapResources(1, &cuda_pbo_resource, 0);

--------------- cuda -------------------------------
texture<uchar, 2, cudaReadModeElementType> tex;
cudaArray *d_imageArray = 0;

__global__ void convertGRBG(uchar4 *d_output, uint width, uint height)
{
    uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
    uint i = __umul24(y, width) + x;

    // input is GR/BG output is BGRA
    if ((x < width) && (y < height)) {

        if ( y & 0x01 ) {
            if ( x & 0x01 ) {  
                d_output[i].x =  (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2;  // B                
                d_output[i].y = (tex2D(tex,x,y));     // G in B
                d_output[i].z = (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2;  // R                    
            } else {
                d_output[i].x = (tex2D(tex,x,y));        //B
                d_output[i].y = (tex2D(tex,x+1,y) + tex2D(tex,x-1,y)+tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/4;  // G
                d_output[i].z = (tex2D(tex,x+1,y+1) + tex2D(tex,x+1,y-1)+tex2D(tex,x-1,y+1)+tex2D(tex,x-1,y-1))/4;   // R
            }
        } else {
            if ( x & 0x01 ) {
                 // odd col = R
                d_output[i].y = (tex2D(tex,x+1,y+1) + tex2D(tex,x+1,y-1)+tex2D(tex,x-1,y+1)+tex2D(tex,x-1,y-1))/4;  // B
                d_output[i].z = (tex2D(tex,x,y));        //R
                d_output[i].y = (tex2D(tex,x+1,y) + tex2D(tex,x-1,y)+tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/4;  // G    
            } else {    
                d_output[i].x = (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2;  // B
                d_output[i].y = (tex2D(tex,x,y));               // G  in R               
                d_output[i].z = (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2;  // R                    
            }
        }                                
    }
}



void initTexture(int imageWidth, int imageHeight, uchar *imagedata)
{

    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsigned);
    cutilSafeCall( cudaMallocArray(&d_imageArray, &channelDesc, imageWidth, imageHeight) ); 
    uint size = imageWidth * imageHeight * sizeof(uchar);
    cutilSafeCall( cudaMemcpyToArray(d_imageArray, 0, 0, imagedata, size, cudaMemcpyHostToDevice) );
    cutFree(imagedata);

    // bind array to texture reference with point sampling
    tex.addressMode[0] = cudaAddressModeClamp;
    tex.addressMode[1] = cudaAddressModeClamp;
    tex.filterMode = cudaFilterModePoint;
    tex.normalized = false; 

    cutilSafeCall( cudaBindTextureToArray(tex, d_imageArray) );
}
0

3 Answers 3

9

There aren't any obvious bugs in your code, but there are several obvious performance opportunities:

1) for best performance, you should use texture to stage into shared memory - see the 'SobelFilter' SDK sample.

2) As written, the code is writing bytes to global memory, which is guaranteed to incur a large performance hit. You can use shared memory to stage results before committing them to global memory.

3) There is a surprisingly big performance advantage to sizing blocks in a way that match the hardware's texture cache attributes. On Tesla-class hardware, the optimal block size for kernels using the same addressing scheme as your kernel is 16x4. (64 threads per block)

For workloads like this, it may be hard to compete with optimized CPU code. SSE2 can do 16 byte-sized operations in a single instruction, and CPUs are clocked about 5 times as fast.

1
  • Thanks - I will take a look at the sobel, I thought I was writing to a texture mapped to a PBO. And yes the demosaic alone isn't worth using the GPU but it's the last step of a lot of other processing and it means I can keep everything local and display without a trip to main memory for demosaicing and returning 4x as many bytes back for display. Nov 10, 2011 at 13:42
1

Based on answer on Nvidia forums, here (for the search engines) is a slightly more optomised version which writes a 2x2 block of pixels in each thread. Although the difference in speed isn't measurable on my setup.

Note it should be called with a gridsize half the size of the image;

dim3 blockSize(16, 16); // for example
dim3 gridSize((width/2) / blockSize.x, (height/2) / blockSize.y);


__global__ void d_convertGRBG(uchar4 *d_output, uint width, uint height)
{
    uint x = 2 * (__umul24(blockIdx.x, blockDim.x) + threadIdx.x);
    uint y = 2 * (__umul24(blockIdx.y, blockDim.y) + threadIdx.y);
    uint i = __umul24(y, width) + x;

    // input is GR/BG output is BGRA
    if ((x < width-1) && (y < height-1)) {
        // x+1, y+1:

        d_output[i+width+1] = make_uchar4( (tex2D(tex,x+2,y+1)+tex2D(tex,x,y+1))/2,  // B                
                                             (tex2D(tex,x+1,y+1)),     // G in B
                                             (tex2D(tex,x+1,y+2)+tex2D(tex,x+1,y))/2,  // R                    
                                             0xff);

        // x, y+1:
        d_output[i+width] =   make_uchar4( (tex2D(tex,x,y+1)),        //B
                                             (tex2D(tex,x+1,y+1) + tex2D(tex,x-1,y+1)+tex2D(tex,x,y+2)+tex2D(tex,x,y))/4,  // G
                                             (tex2D(tex,x+1,y+2) + tex2D(tex,x+1,y)+tex2D(tex,x-1,y+2)+tex2D(tex,x-1,y))/4,   // R
                                             0xff);


        // x+1, y:
        d_output[i+1] =       make_uchar4( (tex2D(tex,x,y-1) + tex2D(tex,x+2,y-1)+tex2D(tex,x,y+1)+tex2D(tex,x+2,y-1))/4,  // B
                                            (tex2D(tex,x+2,y) + tex2D(tex,x,y)+tex2D(tex,x+1,y+1)+tex2D(tex,x+1,y-1))/4,  // G
                                            (tex2D(tex,x+1,y)),        //R
                                            0xff);


        // x, y:
        d_output[i] =         make_uchar4( (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2,  // B
                                             (tex2D(tex,x,y)),               // G  in R           
                                             (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2,  // R                    
                                             0xff);

    }
}
5
  • You don't mention having run this in the CUDA profiler. Finding out what the profiler says about your code should be first think you do when encountering performance issues -- actually, the first think you do even if you don't encounter any performance issues :)
    – Roger Dahl
    Apr 6, 2012 at 2:58
  • @RogerDahl - I'm new to CUDA so running this in the profiler and finding that it's totally dominated by transfer time or by vsnc doesn't help (as I said it was fast enough). Another poster on here saying - Oh you should never use 'foo' because it blocks the pipeline, always use 'bar' - is useful for the future. Apr 6, 2012 at 14:44
  • I'm confused... This was a question about performance, right? You started your question by stating that your CUDA code was running slower than single threaded CPU code, and then went on to ask if you were doing something wrong. The profiler would be the best way for you to find out if you are doing something wrong. Wouldn't the main purpose of all this be to find out what the dominating factor is that is slowing things down? Where did you say that it's already running fast enough, and if it is, what was the purpose of the question?
    – Roger Dahl
    Apr 6, 2012 at 17:04
  • I was using the VS profiler - which said it was very very slow, but that was a user error = the GPU timer reports in 'ms' while the cpu timer reports 's'. But a profiler doesn't tell you that there is a much better technique you haven't tried it just tells you that data i/O dominates Apr 6, 2012 at 17:38
  • I see. I think others may be confused by this questions as well. Maybe you should just delete this and create a new question with your answer.
    – Roger Dahl
    Apr 6, 2012 at 18:19
0

There are many if's and else's in the code. If you structure the code to eliminate all the conditional statements then you will get a huge performance boost as branching is a performance killer. It is indeed possible to remove the branches. There are exactly 30 cases which you will have to code explicitly. I have implemented it on CPU and it does not contain any conditional statements. I am thinking of making a blog explaining it. Will post it once its done.

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.