日志

[原创]简单的CUDA程序：图像二值化处理

已有 2192 次阅读2009-7-28 16:43 |个人分类:高性能计算|系统分类:兴趣爱好| 二值化

常规的C++实现方法（假设图像为R8G8B8格式，每个像素4字节）：

int x, y;
unsigned long offset = 0, p;
unsigned char r, g, b, mi, ma;

for(y = 0; y < HEIGHT; y ++)
for(x = 0; x < WIDTH; x ++)
{
    p = ((unsigned long *)in)[offset];

    b = ((unsigned char *)(&p))[0];
    g = ((unsigned char *)(&p))[1];
    r = ((unsigned char *)(&p))[2];

    mi = min(r, min(g, b));
    ma = max(r, max(g, b));

    if(((unsigned short)ma + (unsigned short)mi) > THRESHOLD * 2)
        out[offset] = 255;
    else
        out[offset] = 0;
    offset ++;
}

用CUDA的实现方法：

1、每个像素一个线程，各像素并行计算，互不干扰。每个线程块256个线程。
2、包含kernel代码和host代码。kernel代码在GPU运行，host代码在CPU运行。
3、kernel运行时，原始图像数据和结果都是保存在显存中。host代码负责分配内存和复制数据。

源代码：

#define THREAD_N    256
#define THRESHOLD    127

__global__ static void binarizeKernel(unsigned char *in, unsigned char *out)
{
    const unsigned long offset = (blockIdx.x * THREAD_N + threadIdx.x);
    unsigned long p = ((unsigned long *)in)[offset];

    unsigned char b = ((unsigned char *)(&p))[0];
    unsigned char g = ((unsigned char *)(&p))[1];
    unsigned char r = ((unsigned char *)(&p))[2];

    unsigned char mi = __min(r, __min(g, b));
    unsigned char ma = __max(r, __max(g, b));

    out[offset] = (((unsigned short)ma + (unsigned short)mi) > THRESHOLD * 2) ? 255 : 0;
}

extern "C" float binarizeCall(unsigned char *in, unsigned char *out, int w, int h)
{
    unsigned char *device_src = 0;
    unsigned char *device_dest = 0;

    cudaMalloc((void **)&device_src, w * h * 3);
    cudaMalloc((void **)&device_dest, w * h);

    unsigned int timer = 0;
    cutCreateTimer(&timer);

    cudaMemcpy(device_src, in, w * h * 3, cudaMemcpyHostToDevice);

    cutStartTimer(timer);

    binarizeKernel<<<w * h / THREAD_N, THREAD_N>>>(device_src, device_dest);
    cudaThreadSynchronize();

    cutStopTimer(timer);

    cudaMemcpy(out, device_dest, w * h, cudaMemcpyDeviceToHost);

    float ms = cutGetTimerValue(timer);
    cutDeleteTimer(timer);

    cudaFree(device_dest);
    cudaFree(device_src);

    return ms;
}