CUDA 图像编程

一个 Grid 分成 按维度分成多个Block，个数为 GridDim.x * GridDim.y
遍历: blockIdx.x , blockIdx.y

一个Block 按维度分成多个Thread，个数为 BlockDim.x * BlockDim.y
Thread 是最小的运行单元
遍历：threadIdx.x , threadIdx.y

图像处理中，一个像素对应到一个thread 中。

从 thread 映射 到 图像 pix 方式 :

  const int ix = blockIdx.x*blockDim.x + threadIdx.x;
  const int iy = blockIdx.y*blockDim.y + threadIdx.y;

务必保证 覆盖所有的 图像 pix 即:

  GridDim.x * blockDim.x > IMG_WIDTH
  GridDim.y * blockDim.y > IMG_HEIGHT

边界没有映射到的 thread 直接返回

if (ix > IMG_WIDTH ||  iy > IMG_HEIGHT) 
	 return 

使用前先定义 Block 块的大小 : BLOCK_W * BLOCK_H 即每个block 中的Thread 个数
为了充分利用线程，不会造成空线程的浪费，Grid 定义为：
（IMG_WIDTH + BLOCK_W - 1) / BLOCK_W ，
(IMG_HEIGHT + BLOCK_H - 1) / BLOCK_H

最后定义 gridSize blockSize 为：
dim3 gridSize((IMG_WIDTH + BLOCK_W - 1) / BLOCK_W, (IMG_HEIGHT + BLOCK_H - 1) / BLOCK_H);
dim3 blockSize(BLOCK_W, BLOCK_H);

实例 ：二维模板滤波

1. kernel 函数，一般在 .cu文件中实现。GPU执行函数

__global__ void k_convolution_2D(	const float *indata,
 									const float *filter,
  									const int width,
 									const int height,
    								const int fitersize_x,
     								const int fitersize_y,
      								float *outdata)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    int halfSize_X = fitersize_x / 2;
    int halfSize_Y = fitersize_y / 2;
    if ((x >= width) || (y >= height))
        return;
    float sum = 0;
    for (int j = -halfSize_Y; j <= halfSize_Y; j++)
    {
        for (int k = -halfSize_X; k <= halfSize_X; k++)
        {
            int real_x = MIRROR(0, (width - 1), (x + k)); // 越界镜像处理
            int real_y = MIRROR(0, (height - 1), (y + j));
            sum += indata[real_y * width + real_x] *
             filter[(j + halfSize_Y * fitersize_x + (k + halfSize_X)];
        }
    }
    outdata[y * width + x] = CLAMP(sum, 0.0, PIXEL_RANGE);
}

2. cpu 调用函数，在CPU 上执行

int convolution_2D(	const float *indata,
 					const float *filter,
 					const int width,
   					const int height,
    				const int fitersize_x,
     				const int fitersize_y,
      				float *outdata)
{
    dim3 gridSize((width + BLOCK_W - 1) / BLOCK_W,
    			 (height + BLOCK_H - 1) / BLOCK_H);
    dim3 blockSize(BLOCK_W, BLOCK_H);
    k_convolution_2D<<<gridSize, blockSize>>>(
    indata, filter, width, height, fitersize_x, fitersize_y, outdata);
    return 0;
}