CUDA和OpenCV实现的图像GAMMA变换

[cpp] view
plaincopy

#include <stdlib.h>

#include <stdio.h>

#include <cv.h>

#include <highgui.h>

#include "cutil_inline.h"

#define GAMMA 0.4

void runTest(int argc, char** argv);

__global__ void testKernel(float* d_idata, float* d_odata, int width, int height, float gamma)

{

unsigned int tid_in_grid_x = blockDim.x*blockIdx.x + threadIdx.x;

unsigned int tid_in_grid_y = blockDim.y*blockIdx.y + threadIdx.y;

unsigned int tid_in_grid = tid_in_grid_y*width + tid_in_grid_x;

d_odata[tid_in_grid] = powf(d_idata[tid_in_grid], gamma);

}

int main(int argc, char** argv)

{

runTest(argc, argv);

CUT_EXIT(argc, argv);

}

void runTest(int argc, char** argv)

{

if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") ) //设备初始化并设定计时器。

cutilDeviceInit(argc, argv);

else

cudaSetDevice(cutGetMaxGflopsDeviceId() );

unsigned int timer = 0;

cutilCheckError(cutCreateTimer(&timer));

cutilCheckError(cutStartTimer(timer));

IplImage* pImg;

if((pImg = cvLoadImage("lena_gray.jpg", CV_LOAD_IMAGE_GRAYSCALE)) != 0 )

{

cvNamedWindow("Image", 1);

cvShowImage("Image", pImg);

unsigned int num_blocks_x = pImg->width/16;

unsigned int num_blocks_y = pImg->height/16;

unsigned int mem_size = sizeof(float)*pImg->widthStep*pImg->height;

float* h_idata = (float*)malloc(mem_size);

float* h_odata = (float*)malloc(mem_size);

float* d_idata;

CUDA_SAFE_CALL(cudaMalloc((void**)&d_idata, mem_size));

float* d_odata;

CUDA_SAFE_CALL(cudaMalloc((void**)&d_odata, mem_size));

for(int i = 0; i < pImg->widthStep*pImg->height; i++) //从图像中复制数据时要特别注意变量类型，char类型的数据计算精度

h_idata[i] = ((uchar)pImg->imageData[i])/255.0; //太低，但是图像中只能存储char类型，所以要进行转换

CUDA_SAFE_CALL(cudaMemcpy(d_idata, h_idata, mem_size, cudaMemcpyHostToDevice));

dim3 grid(num_blocks_x, num_blocks_y, 1);

dim3 threads(16, 16, 1);

testKernel<<<grid, threads>>>(d_idata, d_odata, pImg->width, pImg->height, GAMMA);

CUT_CHECK_ERROR("Kernel execution failed");

CUDA_SAFE_CALL(cudaMemcpy(h_odata, d_odata, mem_size, cudaMemcpyDeviceToHost));

IplImage* oImg = cvCreateImage(cvSize(pImg->width,pImg->height), IPL_DEPTH_8U, 1);

for(int i = 0; i < pImg->widthStep*pImg->height; i++)

oImg->imageData[i] = (uchar)(int)(h_odata[i]*255);

cvNamedWindow("Result", CV_WINDOW_AUTOSIZE);

cvShowImage("Result", oImg);

// cvSaveImage("result.jpg", oImg);

cutilCheckError( cutStopTimer( timer));

printf( "Processing time: %f (ms)/n", cutGetTimerValue( timer));

cutilCheckError( cutDeleteTimer( timer));

cvWaitKey(0);

cvDestroyWindow("Image");

cvReleaseImage(&pImg);

cvDestroyWindow("Result");

cvReleaseImage(&oImg);

free(h_idata);

free(h_odata);

CUDA_SAFE_CALL(cudaFree(d_idata));

CUDA_SAFE_CALL(cudaFree(d_odata));

}

}