使用CUDA查找一堆图像的像素平均值

Question

所以我有一个图像立方体。 512X512X512，我想总结像素明智的图像，并将其保存到最终的结果图像。所以如果所有的像素值为1 ...最终的图像将全部为512.我无法理解索引在CUDA中做到这一点。我认为一个线程的工作就是总结所有512个像素......所以线程总数将是512X512。所以我打算用512块，每块512块。从这里开始，我很难提出如何总结深度的索引。任何帮助将不胜感激。

Answer 1

解决此问题的一种方法是将立方体成像为一组Z片。坐标X，Y指的是图像的宽度和高度，以及Z尺寸中每个幻灯片的Z坐标。每个线程将在Z坐标中迭代以累加值。

考虑到这一点，配置一个内核来启动一个16x16线程块和一个足够块的网格来处理图像的宽度和高度（我假设每像素1个字节的灰度图像）：

#define THREADS 16 
// kernel configuration 
dim3 dimBlock = dim3 (THREADS, THREADS, 1); 
dim3 dimGrid = dim3 (WIDTH/THREADS, HEIGHT/THREADS); 
// call the kernel 
kernel<<<dimGrid, dimBlock>>>(i_data, o_Data, WIDTH, HEIGHT, DEPTH); 

如果你清楚如何索引二维数组，遍历Z尺寸也将是明显的

__global__ void kernel(unsigned char* i_data, unsigned char* o_data, int WIDTH, int HEIGHT, int DEPTH) 
{ 
    // in your kernel map from threadIdx/BlockIdx to pixel position 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    // calculate the global index of a pixel into the image array 
    // this global index is to the first slide of the cube 
    int idx = x + y * WIDTH; 

    // partial results 
    int r = 0; 

    // iterate in the Z dimension 
    for (int z = 0; z < DEPTH; ++z) 
    { 
    // WIDTH * HEIGHT is the offset of one slide 
    int idx_z = z * WIDTH*HEIGHT + idx; 
    r += i_data[ idx_z ]; 
    } 
    // o_data is a 2D array, so you can use the global index idx 
    o_data[ idx ] = r; 
} 

这是一个天真的实现。为了最大化内存吞吐量，数据应该正确对齐。

Answer 2

这可以使用ArrayFire GPU库（免费）轻松完成。在ArrayFire，您可以构建三维阵列类似如下：

两种方法：

// Method 1: 
array data = rand(x,y,z); 
// Just reshaping the array, this is a noop 
data = newdims(data,x*y, z, 1); 

// Sum of pixels 
res = sum(data); 

// Method 2: 
// Use ArrayFire "GFOR" 
array data = rand(x,y,z);res = zeros(z,1); 
gfor(array i, z) { 
    res(ii) = sum(data(:,:,i); 
}