确定CUDA内核的块数和线程数，以及如何使用它们。

Question

我一直在尝试弄清楚如何制作一个简单的内核，以计算二维矩阵中值的平均值，但是我在思考过程中遇到了一些问题。

根据我的deviceQuery输出，我的GPU有16MP，32cores/mp，最大块数为1024x1024x64，每个块最大线程数为1024。

因此，我正在处理一些大型图像。可能是5000px x 3500px或类似的大小。我的一个内核正在对图像中所有像素的某些值取平均值。

现有代码将图像存储为2D数组[rows][cols]。因此，在C中，该内核看起来像您期望的那样，通过循环遍历行和列，并在中间进行计算。

那么，我如何在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];
}
}
}

使用cudaMallocPitch函数声明二维数组的方式与在C语言中使用malloc函数类似，但是没有提到如何在自己的内核中访问该数组。我猜在我的代码中，我将使用cudaMallocPitch调用，然后执行memcpy将数据传输到设备上的二维数组中？

如果有任何提示，请告诉我！谢谢！

- Derek

Answer 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);   
}

Good luck!!!

- user3398689

Answer 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;

将二维数组以这种方式存储到单个连续的内存块中的方法称为行优先存储数据。请参见维基百科文章此处。一旦您将数据布局更改为这种格式，您就可以使用在SDK中显示的缩减，并且您的代码应该会更快，因为您将能够在GPU内核代码中执行更多协作读取操作。

- Samsdram

---

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

Answer 3

以下是我自己代码中的一个简单内核的短代码片段。所有浮点指针都是设备指针。希望这对您有所帮助。

定义和辅助函数：

#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;
}
}

- peakxu

---

如果我理解iDivUP的作用，你可以通过整数截断简化逻辑：返回(a+b-1)/b; - Erich Mirabal