Использование 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 Библиотека графического процессора (бесплатно). В 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);
}