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
Post a Comment