取平均超过块 Cuda答案

【问题标题】：Take average over blocks Cuda取平均超过块 Cuda
【发布时间】：2013-06-25 22:33:38
【问题描述】：

Cuda 说共享内存只能被同一个块中的数据共享。但是一个块最多只能有 1024 个线程。如果我有一个巨大的矩阵，并且想用最大化的线程取它们的平均值怎么办。

以此为例。（我没有在一个block中使用最大化线程，只是作为一个demo）

#include <iostream>
#include <stdio.h>

__global__ void
kernel(int *a, int dimx, int dimy)
{
int ix = blockDim.x * blockIdx.x + threadIdx.x;
int iy = blockDim.y * blockIdx.y + threadIdx.y;

int idx = iy * dimx + ix;

__shared__ int array[64];

a[idx] = a[idx] + 1;

array[idx] = a[idx];

__syncthreads();

int sum=0;
for(int i=0; i<dimx*dimy; i++)
{
    sum += array[i];
}

int average = sum/(dimx*dimy+1.0f);

a[idx] = average;

 }

int
main()
{
int dimx = 8;
int dimy = 8;
int num_bytes = dimx*dimy*sizeof(int);

int *d_a=0, *h_a=0; // device and host pointers
h_a = (int*)malloc(num_bytes);

for (int i=0; i < dimx*dimy; i++){
    *(h_a+i) = i;
}
cudaMalloc( (void**)&d_a, num_bytes );

//cudaMemset( d_a, 0, num_bytes );

cudaMemcpy( d_a, h_a, num_bytes, cudaMemcpyHostToDevice);

dim3 grid, block;
block.x = 4;
block.y = 4;
grid.x = dimx / block.x;
grid.y = dimy / block.y;

kernel<<<grid, block>>>(d_a, dimx, dimy);

cudaMemcpy( h_a, d_a, num_bytes, cudaMemcpyDeviceToHost );


std::cout << "the array a is:" << std::endl;
for (int row = 0; row < dimy; row++)
{
    for (int col =0; col < dimx; col++)
    {
        std::cout << h_a[row * dimx + col] << " ";
    }
    std::cout << std::endl;
}

free(h_a);
cudaFree(d_a);
}

我创建了四个块，并希望结果是所有块的平均值。现在的结果是：

the array a is:
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
3 3 3 3 4 4 4 4 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12 
11 11 11 11 12 12 12 12

每个块都有自己的平均值，而不是整体平均值。我怎样才能取所有块的平均值？

我是 Cuda 的新手。欢迎任何相关的答案。

【问题讨论】：

查看 CUDA SDK 中的一个缩减示例，看看这种事情是如何完成的。

标签： cuda gpgpu

【解决方案1】：

最简单的方法是启动多个内核，这样您就可以计算每个块的平均值，将它们写入全局内存，然后启动另一个内核来处理前一个内核的每个块的结果。根据您的数据维度，您可能需要多次重复此操作。

例如（伪代码）

template <typename T>
__global__ reduce(T* data, T* block_avgs)
{
    //find the per-block average, write it out to block_avgs
    //...
}

//in your caller:
loop while you have more than 1 block:
    call kernel using result from prev. iteration
    update grid_dim and block_dim

这是必要的，因为 CUDA 中没有块间同步。您的问题是减少的非常简单的应用。查看并行缩减示例at the nvidia samples page 以更好地了解缩减。

【讨论】：