弄清楚一个 cuda 内核有多少块和线程，以及如何使用它们

Posted 2023-03-23

【中文标题】弄清楚一个 cuda 内核有多少块和线程，以及如何使用它们

【英文标题】：figuring out how many blocks and threads for a cuda kernel, and how to use them

【发布时间】：2011-06-15 06:03:42

【问题描述】：

我一直在试图弄清楚如何制作我认为简单的内核来取 2d 矩阵中的值的平均值，但是我在直接处理我的思考过程时遇到了一些问题。

根据我的 deviceQuery 输出，我的 GPU 有 16MP，32cores/mp，blocks max 是 1024x1024x64，我有一个 max threads/block=1024。

所以，我正在处理一些大图像。也许是 5000 像素 x 3500 像素或类似的东西。我的内核之一是对图像中所有像素的一些值取平均值。

现有代码将图像存储为二维数组 [rows][cols]。因此，在 C 语言中，内核看起来就像您所期望的那样，在行上循环，在 cols 上循环，计算在中间。

那么如何在 CUDA 中设置这段代码的维度计算部分呢？我已经查看了 SDK 中的缩减代码，但这是针对一维数组的。它没有提到当你有一些 2D 时如何设置块和线程的数量。

我想我实际上需要像这样设置它，这就是我希望有人插话并提供帮助的地方：

num_threads=1024;
blocksX = num_cols/sqrt(num_threads);
blocksY = num_rows/sqrt(num_threads);
num_blocks = (num_rows*num_cols)/(blocksX*blocksY);

dim3 dimBlock(blocksX, blocksY, 1);
dim3 dimGrid(num_blocks, 1, 1);

这似乎对设置有意义吗？

然后在内核中，要处理特定的行或列，我必须使用

rowidx = (blockIdx.x*blockDim.x)+threadId.x colidx = (blockIdx.y*blockDim.y)+threadId.y

至少我认为这对于获取行和列是有效的。

然后我将如何访问内核中的特定行 r 和列 c？在 cuda 编程指南中，我找到了以下代码：

// Host code int width = 64, height = 64;
float* devPtr; size_t pitch;
cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height);
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
// Device code __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height)

for (int r = 0; r < height; ++r)

float* row = (float*)((char*)devPtr + r * pitch);
for (int c = 0; c < width; ++c)

float element = row[c];

这看起来类似于您在 C 中使用 malloc 声明二维数组的方式，但它没有提到在您自己的内核中访问该数组。我想在我的代码中，我将使用那个 cudaMallocPitch 调用，然后执行 memcpy 将我的数据放入设备上的二维数组？

感谢任何提示！谢谢！

【参考方案1】：

最近，我以以下方式思考了这个问题。

// Grid and block size
const dim3 blockSize(16,16,1);
const dim3 gridSize(numRows, numCols, 1); 
// kernel call
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols

gridsize = 块数 blocksize = 每个块的线程数

这里是对应的内核

__global__ void rgba_to_greyscale(const uchar4* const rgbaImage,
                       unsigned char* const greyImage,
                       int numRows, int numCols)
 
    int idx = blockIdx.x + blockIdx.y * numRows;
    uchar4 pixel     = rgbaImage[idx]; 
    float  intensity = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z;  
    greyImage[idx]   = static_cast<unsigned char>(intensity);   

祝你好运！！！

【参考方案2】：

对于这样的性能应用，您需要将 2D 矩阵信息作为单个数组存储在内存中。因此，如果您有一个 M x N 矩阵，则可以将其存储在长度为 M*N 的单个数组中。

所以如果你想存储 2x2 矩阵

(1 , 2)
(3 , 4)

然后您创建一个数组，并使用以下方法初始化第 i 行和第 j 列的元素。

int rows=2;
int cols=2;
float* matrix = malloc(sizeof(float)*rows*cols);
matrix[i*cols+j]=yourValue;
//element 0,0
matrix[0*cols+0]=1.0;
//element 0,1
matrix[0*cols+1]=2.0;
//element 1,0
matrix[1*cols+0]=3.0;
//element 1,1
matrix[1*cols+1]=4.0;

这种获取二维数组并以这种方式将其存储为单个连续内存块的方式称为以行优先顺序存储数据。参见***文章here。一旦您将数据的布局更改为这种格式，您就可以使用 SDK 中显示的缩减，并且您的代码应该会更快，因为您将能够在 GPU 内核代码中进行更多的合并读取。

【讨论】：

我同意这是解决此问题的最简单（并且可能是最有效）的方法。我唯一关心的是精度：如果您正在对具有高精度像素的非常大的图像进行求和，那么您可能会用完位，因此请确保使用足够大的数据类型。或者，您可以修改减少以计算运行平均值而不是总和。

【参考方案3】：

下面是一个简短的 sn-p，带有我自己的代码中的简单内核。浮点指针都是设备指针。希望这会有所帮助。

定义和帮助函数：

#define BLOCK_SIZE 16

int iDivUp(int a, int b)
    return (a % b != 0) ? (a / b + 1) : (a / b);

块大小计算：

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
dim3 dimGridProj(iDivUp(width,BLOCK_SIZE), iDivUp(height,BLOCK_SIZE));

主机调用：

calc_residual<<<dimGridProj, dimBlock>>>(d_image1, d_proj1, d_raynorm1, d_resid1, width, height);

内核：

__global__ void calc_residual(float *d_imagep, float *d_projp, float *d_raysump, float *d_residualp, int width, int height)

    int iy = blockDim.y * blockIdx.y + threadIdx.y;
if (iy >= height) 
    return;

int ix = blockDim.x * blockIdx.x + threadIdx.x;
if (ix >= width) 
    return;

int idx = iy * width + ix;
float raysumv = d_raysump[idx];
if (raysumv > 0.001) 
    d_residualp[idx] = (d_projp[idx]-d_imagep[idx])/raysumv;
 
else
    d_residualp[idx] = 0;

【讨论】：

如果我理解 iDivUP 在做什么，您可以通过整数截断简化逻辑：return (a+b-1)/b;