CUDA — реализация медианного фильтра не дает желаемых результатов

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

Вам может повезти со следующим:

$ cat t376.cu
#include <stdlib.h>
#include <stdio.h>

#define DCOLS 1024
#define DROWS 256

typedef struct {
size_t step;
size_t rows;
size_t cols;
unsigned char *data;
} mat;// define the threads and grids for CUDA
#define BLOCK_ROWS 32
#define BLOCK_COLS 16

// define kernel dimensions
#define MEDIAN_LENGTH 9

// this is the error checking part for CUDA
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
if (code != cudaSuccess)
{
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}__global__ void FilterKernel (unsigned char *d_input_img, unsigned char *d_output_img, int d_iRows, int d_iCols)

{

unsigned int row = blockIdx.y*blockDim.y + threadIdx.y;
unsigned int col = blockIdx.x*blockDim.x + threadIdx.x;
unsigned char window[MEDIAN_LENGTH];

if(col>=d_iCols || row>=d_iRows)
return;

window[0]= (row==0||col==0) ? 0 :                 d_input_img[(row-1)*d_iCols+(col-1)];
window[1]= (row==0) ? 0 :                         d_input_img[(row-1)*d_iCols+col];
window[2]= (row==0||col==d_iCols-1) ? 0 :         d_input_img[(row-1)*d_iCols+(col+1)];
window[3]= (col==0) ? 0 :                         d_input_img[row*d_iCols+(col-1)];
window[4]=                                        d_input_img[row*d_iCols+col];
window[5]= (col==d_iCols-1) ? 0 :                 d_input_img[row*d_iCols+(col+1)];
window[6]= (row==d_iRows-1||col==0) ? 0 :         d_input_img[(row+1)*d_iCols+(col-1)];
window[7]= (row==d_iRows-1) ? 0 :                 d_input_img[(row+1)*d_iCols+col];
window[8]= (row==d_iRows-1||col==d_iCols-1) ? 0 : d_input_img[(row+1)*d_iCols+(col+1)];

// Order elements
for (unsigned int j=0; j<5; ++j)
{
// Find position of minimum element
unsigned char temp = window[j];
unsigned int  idx  = j;
for (unsigned int l=j+1; l<9; ++l)
if (window[l] < temp){ idx=l; temp = window[l];}
// Put found minimum element in its place
window[idx] = window[j];
window[j] = temp;
}

d_output_img[row*d_iCols + col] = (window[4]);

}

void take_input(const mat& input, const mat& output)
{

unsigned char *device_input;
unsigned char *device_output;

size_t d_ipimgSize = input.step * input.rows;
size_t d_opimgSize = output.step * output.rows;

gpuErrchk( cudaMalloc( (void**) &device_input, d_ipimgSize) );
gpuErrchk( cudaMalloc( (void**) &device_output, d_opimgSize) );

gpuErrchk( cudaMemcpy(device_input, input.data, d_ipimgSize, cudaMemcpyHostToDevice) );

dim3 Threads(BLOCK_COLS, BLOCK_ROWS);  // 512 threads per block
dim3 Blocks((input.cols + Threads.x - 1)/Threads.x, (input.rows + Threads.y - 1)/Threads.y);

//int check = (input.cols + Threads.x - 1)/Threads.x;
//printf( "blockx %d", check);

FilterKernel <<< Blocks, Threads >>> (device_input, device_output, input.rows, input.cols);
gpuErrchk(cudaDeviceSynchronize());
gpuErrchk(cudaGetLastError());

gpuErrchk( cudaMemcpy(output.data, device_output, d_opimgSize, cudaMemcpyDeviceToHost) );

//printf( "num_rows_cuda %d", num_rows);
//printf("\n");

gpuErrchk(cudaFree(device_input));
gpuErrchk(cudaFree(device_output));

}

int main(){
mat input_im, output_im;
input_im.rows  = DROWS;
input_im.cols  = DCOLS;
input_im.step  = input_im.cols;
input_im.data  = (unsigned char *)malloc(input_im.step*input_im.rows);
output_im.rows = DROWS;
output_im.cols = DCOLS;
output_im.step = input_im.cols;
output_im.data = (unsigned char *)malloc(output_im.step*output_im.rows);

for (int i = 0; i < DCOLS*DROWS; i++) {
output_im.data[i] = 0;
input_im.data[i] = 0;
int temp = (i%DCOLS);
if (temp == 5) input_im.data[i] = 20;
if ((temp > 5) && (temp < 15)) input_im.data[i] = 40;
if (temp == 15) input_im.data[i] = 20;
}

take_input(input_im, output_im);
for (int i = 2*DCOLS; i < DCOLS*(DROWS-2); i++)
if (input_im.data[i] != output_im.data[i]) {printf("mismatch at %d, input: %d, output: %d\n", i, (int)input_im.data[i], (int)output_im.data[i]); return 1;}
printf("Success\n");
return 0;
}$ nvcc -o t376 t376.cu
$ ./t376
Success
$

несколько заметок:

1. Я проверил это (не используя OpenCV) для простого случая, который я вставил в код.
2. Ваше использование window было излишне сложно. Обратите внимание, что так, как у вас есть, каждая нить будет создавать свою собственную локальную копию window не зависит и невидим для других потоков. (возможно, вы намеревались использовать общую память здесь? Но я отвлекся.)
3. Ваша сортировка была нарушена. Я изменил его до версии, которая, я думаю, будет работать.
4. заменены типы данных пикселей на unsigned char на протяжении
5. x а также y были в замешательстве, поэтому я изменил их на row а также col который кажется менее запутанным.
6. немного улучшена проверка ошибок ядра
7. Есть много способов оптимизировать это. Однако ваша разумная цель состояла в том, чтобы сначала заставить что-то работать правильно. Поэтому я не буду тратить много времени на оптимизацию, за исключением того, чтобы указать на две основные области совместной памяти для повторного использования window данные и улучшенная процедура сортировки.
8. Вам нужно будет изменить это соответствующим образом для openCV
9. Обратите внимание, что если вы измените его на DROWS = 1024 и DCOLS = 256, он все равно будет работать.

РЕДАКТИРОВАТЬ: после прочтения ваших комментариев о том, что что-то все еще не работает, кажется, что ваш код OpenCV не настроен должным образом для подачи на одноканальное 8-битное изображение в оттенках серого (CV_8UC1) в а также от ваш take_input функция. Проблема возникает из этой строки:

cv::Mat input = cv::imread("C:/Users/OCT/Documents/Visual Studio 2008/Projects/MedianFilter/MedianFilter/pic1.bmp",1);

1 параметр передается imread определяет загрузку изображения RGB. Ссылаться на непрочитанная документация:

Now we call the imread function which loads the image name specified by the first argument (argv[1]). The second argument specifies the format in what we want the image. This may be:

CV_LOAD_IMAGE_UNCHANGED (<0) loads the image as is (including the alpha channel if present)
CV_LOAD_IMAGE_GRAYSCALE ( 0) loads the image as an intensity one
CV_LOAD_IMAGE_COLOR (>0) loads the image in the RGB format

Возможно, вам повезет больше, если вы укажете CV_LOAD_IMAGE_GRAYSCALE там вместо 1,

Или же вы должны изучить, как загрузить изображение, чтобы оно оказалось CV_8UC1 тип.

Но если вы передадите это input в take_input как есть, это точно не сработает.