как сделать ядро ​​гистограммы CUDA?

Question

Я пишу ядро ​​CUDA для гистограммы на картинке, но я понятия не имел, как вернуть массив из ядра, и этот массив изменится, когда другой поток прочитает его.Любое возможное решение для этого?

__global__ void Hist(
    TColor *dst, //input image
    int imageW,
    int imageH,
 int*data
){
    const int ix = blockDim.x * blockIdx.x + threadIdx.x;
    const int iy = blockDim.y * blockIdx.y + threadIdx.y;

if(ix < imageW && iy < imageH)
{
  int pixel = get_red(dst[imageW * (iy) + (ix)]);
                  //this assign specific RED value of image to pixel

  data[pixel] ++; // ?? problem statement ...
 }
}

@ para d_dst: входное изображение TColor равно float4.

@ para data: массив для размера гистограммы [255]

extern "C" void
cuda_Hist(TColor *d_dst, int imageW, int imageH,int* data) 
{
  dim3 threads(BLOCKDIM_X, BLOCKDIM_Y);
  dim3 grid(iDivUp(imageW, BLOCKDIM_X), iDivUp(imageH, BLOCKDIM_Y));
  Hist<<<grid, threads>>>(d_dst, imageW, imageH, data);
}

Answer 1

Вы смотрели образец SDK? Пример «гистограммы» доступен в CUDA SDK (в настоящее время версия 3.0 на сайте разработчиков NVIDIA , бета-версия 3.1 доступна для зарегистрированных разработчиков ).

Документация с примером хорошо объясняет, как обрабатывать суммирование, используя глобальную атомизацию памяти на GPU или собирая результаты для каждого блока в отдельности, а затем делая отдельное сокращение (на хосте или GPU).

Answer 2

Гистограмма не особенно эффективна при реализации с CUDA (или с GPGPU в целом) - обычно вам нужно генерировать много частичных гистограмм в разделяемой памяти и затем суммировать их. Возможно, вы захотите сохранить эту конкретную задачу на процессоре.

Answer 3

Вам придется либо использовать атомарную функцию, чтобы заблокировать другой поток от использования той же памяти, или использовать частичную гистограмму. В любом случае, это не так эффективно, если только входное изображение не очень большое.