NVIDIA CUDA简单的CUDA程序：图像二值化处理

 

常规的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;
}
 

 

运行速度比较（CPU用QX6600 2.4GHz单线程，GPU用Quadro FX5800，GT200核心）

图像大小1280×800（1M像素）：CPU耗时5.5ms，GPU耗时0.16ms。
图像大小3264×2448（8M像素）：CPU耗时41.2ms，GPU耗时0.83ms。

用ION平台的ATOM 230 CPU和集成GPU比较：

图像大小1280×800（1M像素）：CPU耗时17.3ms，GPU耗时1.8~2.2ms。

结论：

1、GT200在做高度并行的图形处理运算时，速度可以达到2.4GHz CPU单线程时的30～50倍，即使是ION平台的集成GPU也比普通的CPU单线程时要快。

2、ION的ATOM CPU性能太差。