遍历图像设计threadIdx blockIdx blockDim 解释

最近进行gpu编程，受困于杂事，对threadIdx blockIdx blockDim理解存在偏颇导致浪费不少时日。遂整理资料加深认识。

以下是最常见的gpu图像遍历方式：

dim3 dimBlock( 线程数1, 线程数2,1) ;//最开始，望文生义而把后面数据理解为 threadIdx 的值导致大量时间浪费在这上面。这边就是线程数，不是块数。不要被网上代码迷惑

dim3 dimGrid( ceil( ( 图像宽 + 线程数1 - 1)/线程数1 ) ,ceil( ( 图像高 + 线程数 2-1) /线程数2) , 1 ) ;

Cal<<<dimGrid , dimBlock>>>(图像数据首地址 …… );

__global__ void Cal(unsigned char * srcData ……)

{
	const int idx = threadIdx.x + blockIdx.x * blockDim.x;
	const int idy = threadIdx.y + blockIdx.y * blockDim.y;
…………
         if(idx  < 图像宽 && idy < 图像高) //图像宽 ，图像高通过参数传入
	{

<span style="white-space: pre;">		</span>……//对图像进行相应单像素操作

	}
}

按照上面访问方式 就可以进行图像逐像素遍历。 各个元素解释如下：

dim3             ：  定义如下

  struct __device_builtin__ dim3
{
    unsigned int x, y, z;
#if defined(__cplusplus)
    __host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
    __host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
    __host__ __device__ operator uint3(void) { uint3 t; t.x = x; t.y = y; t.z = z; return t; }
#endif /* __cplusplus */
};

blockDim.x ： 块操作x方向维数 ， 对应   ceil( ( 图像宽 + 线程数1 - 1)/线程数1 )  所在位置指定值。

blockDim.y ： 快操作y方向维数 ， 对应   ceil( ( 图像高 + 线程数 2-1) /线程数2)   所在位置指定值。

threadIdx.x ： 当前线程 x 方向值 ， 范围 0~ 线程数1 -1 。

threadIdx.y ： 当前线程y方向值   ， 范围 0~ 线程数2 - 1 。

blockIdx.x   ： 当前块 x方向值 ，       范围 0~ceil( ( 图像宽 + 线程数1 - 1)/线程数1 )  -  1

blockIdx.y   ： 当前块 y方向值 ，       范围 0~ceil( ( 图像高 + 线程数 2-1) /线程数2) - 1

通过以上 知识，可以计算得到   idx 的最小值= 0 ， idx最大值<=  图像宽 + 线程数1 - 1 ，并且遍历 最大最小值之间所有值，因此可以把宽度方向遍历完毕。同理，y方向也刚好每行遍历一次，不存在重复遍历元素情况。

在之后，因为有共有参数，嫌这样方式慢，我更改为二维方式访问。频繁出错，出错原因就在于  忽略  “遍历 最大最小值之间所有值” 这个条件导致出错。

最后遍历方式：

//块，线程设置

dim3 dimBlock(线程数1, 线程数2,1) ;

dim3 dimGrid(1, 块数2, 1 ) ;

//调用

Cal<<<dimGrid , dimBlock>>>(图像数据首地址 …… );

__global__ void Cal(unsigned char * srcData ……)

{

const int idx = threadIdx.x ;

const int idy = threadIdx.y + blockIdx.y * blockDim.y ;

//载入共享数据

……

for(int dy = idy + offsetLoc;dy < 图像高; dy += 线程数2*blockDim.y)

{

for(int dx = idx ; dx < 图像宽; dx += 线程数1)

{

//逐像素操作

}

}

}

之所以按照上面访问方式，是因为有使用共享内存的需要。并不是最优方案，按照自己需求更改访问方式。