第一步获取GPU硬件有多少block,有多少线程,代码如下
int getThreadNum()
{
cudaDeviceProp prop;
int count;
HANDLE_ERROR(cudaGetDeviceCount(&count));
printf("gpu num %d\n", count);
HANDLE_ERROR(cudaGetDeviceProperties(&prop, 0));
printf("max thread num: %d\n", prop.maxThreadsPerBlock);
printf("max grid dimensions: %d, %d, %d)\n",
prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
return prop.maxThreadsPerBlock;
}
我的机器block数量为(2147483647,65535,65535),每个block线程数量为1024
第二步根据实际矩阵大小看需要多少block
例如有一个 1920×1080的一维矩阵,应该如何设计blockNum和ThreadNum ?
int threadNum = getThreadNum();
int blockNum = (width * height - 0.5) / threadNum + 1;
我的机器block数量非常多,所以用1维就够了,ThreadNum用机器最大的线程1024即可,那么blockNum就等于总数除以每个block的线程数,为什么要 -0.5和+1是为了防止整除、进一位等,总之加上为好
第三步开始使用
conv<<<blockNum, threadNum>> >
(imgGpu, kernelGpu, resultGpu, width, height, kernelSize);
__global__ void conv(float *img, float *kernel, float *result,
int width, int height, int kernelSize)
{
int ti = threadIdx.x;
int bi = blockIdx.x;
int id = (bi * blockDim.x + ti);
if(id >= width * height)
{
return;
}
int row = id / width;
int col = id % width;
for(int i = 0; i < kernelSize; ++i)
{
for(int j = 0; j < kernelSize; ++j)
{
float imgValue = 0;
int curRow = row - kernelSize / 2 + i;
int curCol = col - kernelSize / 2 + j;
if(curRow < 0 || curCol < 0 || curRow >= height || curCol >= width)
{}
else
{
imgValue = img[curRow * width + curCol];
}
result[id] += kernel[i * kernelSize + j] * imgValue;
}
}
}