data type conversion - Cuda issue on converting image to grayscale -


i having issue following code. following code takes input image , should save grayscale of it. unfortunately, seems perform expected behavior processing part of image , not whole. seems problems occurs in cudamemcpy device host.

i believe got issue while allocating memory in cuda.

__global__ void rgb2graycudakernel(unsigned char *inputimage, unsigned char *grayimage, const int width, const int height)  {       int ty = (blockidx.x * blockdim.x) + threadidx.x;     //int tx = (blockidx.x * blockdim.x) + threadidx.x;     int tx = (blockidx.y * blockdim.y) + threadidx.y;      if( (ty < height && tx<width) )      {              float graypix = 0.0f;             float r = static_cast< float >(inputimage[(ty * width) + tx]);                       float g = static_cast< float >(inputimage[(width * height) + (ty * width) + tx]);             float b = static_cast< float >(inputimage[(2 * width * height) + (ty * width) + tx]);              graypix = (0.3f * r) + (0.59f * g) + (0.11f * b);               grayimage[(ty * width) + tx] = static_cast< unsigned char >(graypix);         }    }  //***************************************rgb2gray function, call of kernel in here ************************************* void rgb2graycuda(unsigned char *inputimage, unsigned char *grayimage, const int width, const int height) {     unsigned char *inputimage_c, *grayimage_c;     const int sizee= (width*height);      // **********memory allocation pointers , cuda******************       cudamalloc((void **) &inputimage_c, sizee);     checkcudaerror("im not alloc!");     cudamalloc((void **) &grayimage_c, sizee);     checkcudaerror("gray not alloc !");  //***********copy device*************************     cudamemcpy(inputimage_c, inputimage, sizee*sizeof(unsigned char), cudamemcpyhosttodevice);     checkcudaerror("im not send !");     cudamemcpy(grayimage_c, grayimage, sizee*sizeof(unsigned char), cudamemcpyhosttodevice);     checkcudaerror("gray not send !");     dim3 thrb(32,32);     dim3 numb (ceil(width*height/1024)); //**************execute kernel (timer in here)**************************     nstimer kerneltime = nstimer("kerneltime", false, false);     kerneltime.start();      rgb2graycudakernel<<<numb,1024>>> (inputimage_c, grayimage_c, width, height);     checkcudaerror("kernel!");     kerneltime.stop(); //**************copy host*************************     printf("/c");     cudamemcpy(grayimage, grayimage_c, sizee*sizeof(unsigned char), cudamemcpydevicetohost);     checkcudaerror("receiving data cpu failed!");  //*********************free memory***************************     cudafree(inputimage_c);     cudafree(grayimage_c);   //**********************print time****************   cout << fixed << setprecision(6); cout << "rgb2gray (cpu): \t\t" << kerneltime.getelapsed() << " seconds." << endl;  } 

const int sizee= (width*height);  

should be:

const int sizee= (width*height*3);  

for rgb data (1 byte per channel).

i believe in bitmap images, colors should interleaved in:

rgb of pixel1, rgb of pixel 2 ... rgb of pixel width*height 

therefore kernel should be:

__global__ void rgb2graycudakernel(unsigned char *inputimage, unsigned char *grayimage, const int width, const int height)  {       int tx = (blockidx.y * blockdim.y) + threadidx.y;     int ty = (blockidx.x * blockdim.x) + threadidx.x;      if( (ty < height && tx<width) )      {             unsigned int pixel = ty*width+tx;             float graypix = 0.0f;             float r = static_cast< float >(inputimage[pixel*3]);                       float g = static_cast< float >(inputimage[pixel*3+1]);             float b = static_cast< float >(inputimage[pixel*3+2]);              graypix = (0.3f * r) + (0.59f * g) + (0.11f * b);               grayimage[pixel] = static_cast< unsigned char >(graypix);         }    } 

also, saw luminosity calculated 0.21 r + 0.72 g + 0.07 b.


Comments

Popular posts from this blog

Email notification in google apps script -

c++ - Difference between pre and post decrement in recursive function argument -

javascript - IE11 incompatibility with jQuery's 'readonly'? -