二维阵列中的Cuda约简

Question

我想计算Cuda中整个图像的值的平均值。为了测试2D阵列中的归约是如何工作的，我在下面编写了这个内核。最终输出o应该是所有图像值的总和。输入g是在每个像素中具有值1的2D阵列。但是这个程序的结果是0作为总和。我觉得有点奇怪。

我模仿本教程中1D阵列的缩小http://developer.download.nvidia.com/compute/cuda/1.1-Beta/x86_website/projects/reduction/doc/reduction.pdf我写这个二维表格。我是库达的新手。欢迎对潜在的错误和改进提出建议！

只需添加一条评论。我知道计算1D数组中的平均值是有意义的。但我想开发更多，测试更复杂的还原行为。这可能不对。但这只是一个测试。希望有人能给我更多关于减少常见做法的建议。

#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
cudaEvent_t start, stop;
float elapsedTime;
__global__ void 
reduce(float *g, float *o, const int dimx, const int dimy)
{
extern __shared__ float sdata[];
unsigned int tid_x = threadIdx.x;
unsigned int tid_y = threadIdx.y;
unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
unsigned int j = blockDim.y * blockIdx.y + threadIdx.y; 
if (i >= dimx || j >= dimy)
    return;
sdata[tid_x*blockDim.y + tid_y] = g[i*dimy + j];
__syncthreads();
for(unsigned int s_y = blockDim.y/2; s_y > 0; s_y >>= 1)
{
    if (tid_y < s_y)
    {
        sdata[tid_x * dimy + tid_y] += sdata[tid_x * dimy + tid_y + s_y];
    }
    __syncthreads();
}
for(unsigned int s_x = blockDim.x/2; s_x > 0; s_x >>= 1 )
{
    if(tid_x < s_x)
    {
        sdata[tid_x * dimy] += sdata[(tid_x + s_x) * dimy];
    }
    __syncthreads();
}
float sum;
if( tid_x == 0 && tid_y == 0)
{ 
    sum = sdata[0];
    atomicAdd (o, sum);   // The result should be the sum of all pixel values. But the program produce 0
}
//if(tid_x==0 && tid__y == 0 ) 
    //o[blockIdx.x] = sdata[0];
}
int
main()
{   
int dimx = 320;
int dimy = 160;
int num_bytes = dimx*dimy*sizeof(float);
float *d_a, *h_a, // device and host pointers
            *d_o=0, *h_o=0;
h_a = (float*)malloc(num_bytes);
h_o = (float*)malloc(sizeof(float));
srand(time(NULL));

for (int i=0; i < dimx; i++)
{   
    for (int j=0; j < dimy; j++)
    {
        h_a[i*dimy + j] = 1;
    }
}
cudaMalloc( (void**)&d_a, num_bytes );
cudaMalloc( (void**)&d_o, sizeof(int) );
cudaMemcpy( d_a, h_a, num_bytes, cudaMemcpyHostToDevice);
cudaMemcpy( d_o, h_o, sizeof(int), cudaMemcpyHostToDevice); 
dim3 grid, block;
block.x = 4;
block.y = 4;
grid.x = dimx / block.x;
grid.y = dimy / block.y;
cudaEventCreate(&start);
cudaEventRecord(start, 0);
int sizeofSharedMemory = dimx*dimy*sizeof(float);
reduce<<<grid, block, sizeofSharedMemory>>> (d_a, d_o, block.x, block.y);
cudaEventCreate(&stop);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
std::cout << "This kernel runs: " << elapsedTime << "ms" << std::endl; 
std::cout << block.x << " " << block.y << std::endl;
std::cout << grid.x << " " << grid.y << std::endl;
std::cout << dimx <<  " " << dimy << " " << dimx*dimy << std::endl;
cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );
std::cout << "The sum is:" << *h_o << std::endl;
free(h_a);
free(h_o);
cudaFree(d_a);
cudaFree(d_o);
}

Answer 1

如果进行基本的cuda错误检查，您会发现reduce内核甚至没有运行。原因如下：

int dimx = 320;
int dimy = 160;
...
int sizeofSharedMemory = dimx*dimy*sizeof(float); // = 204800
reduce<<<grid, block, sizeofSharedMemory>>> (d_a, d_o, block.x, block.y);
                          ^
                          |
                         204800 is illegal here

您不能动态（或以任何其他方式）请求204800字节的共享内存。最大值略小于48K字节。

如果你做了正确的cuda错误检查，你会发现你的内核没有运行，并且会得到一条指导性的错误消息，表明启动配置（<<<…>>之间的数字）无效。共享内存是在每个块的基础上请求的，当每个块仅由4x4线程数组组成时，您需要请求足够的共享内存来覆盖整个2D数据集可能是不明智的。您可能只需要足够的数据就可以访问每个4x4线程数组。

在使用cuda错误检查正确地插入代码并检测并更正所有错误之后，然后使用cuda-memcheck运行代码。这将进行额外级别的错误检查，以指出任何内核访问错误。如果出现未指定的启动失败，也可以使用cuda-memcheck，这可能有助于确定问题所在。

完成这些基本的故障排除步骤后，向他人寻求帮助可能是有意义的。但是，首先要使用你所得到的工具的力量。

我还想指出另一个错误，然后你再回来发布这段代码，寻求帮助。

这将毫无用处：

std::cout << "The sum is:" << *h_o << std::endl;
cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );

在将总和从设备复制到主机之前，您正在打印出总和。颠倒这些步骤的顺序：

cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );
cudaMemcpy( h_o, d_o, sizeof(int), cudaMemcpyDeviceToHost );
std::cout << "The sum is:" << *h_o << std::endl;