Я пытался реализовать алгоритм для медианного фильтра, представленного в статье Wiki:http://en.wikipedia.org/wiki/Median_filter#2D_median_filter_pseudo_code

Насколько я знаю, я знаю, что то, что я реализовал, правильно. Однако, когда я просматриваю результаты, я не могу получить вывод, который похож на вывод, произведенныйmedian blur функция в OpenCV. В настоящий момент меня не волнует ускорение моего кода за счет использования общей памяти или памяти текстур. Я просто хотел бы, чтобы все заработало в первую очередь. Размер моего входного изображения1024 x 256 пиксели.

Что я делаю неправильно? Есть ли утечка потоков в моем коде? Я знаю, что я должен использовать разделяемую память для предотвращения глобального чтения, потому что в настоящее время я много читаю данные из глобальной памяти.

http://snag.gy/OkXzP.jpg - первое изображение является входным, второе изображение является результатом моего алгоритма, а третье является openCVmedianblur результат функции. В идеале я хотел бы, чтобы мой алгоритм выводил тот же результат, что иmedianblur функция.

Это весь код, который я написал:

реализация ядра

#include "cuda.h"
#include "cuda_runtime_api.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include "highgui.h"
//#include "opencv2/core/imgproc.hpp"
//#include "opencv2/core/gpu.hpp"
#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>

// includes, project

#include "cufft.h"
#include "cublas_v2.h"
#include "CUDA_wrapper.h"   // contains only func_prototype for function take_input()


// define the threads and grids for CUDA
#define BLOCK_ROWS 32
#define BLOCK_COLS 16

// define kernel dimensions
#define KERNEL_DIMENSION 3
#define MEDIAN_DIMENSION 3
#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);
   }
}

// create two vars for the rows and cols of the image
    int d_imgRows;
    int d_imgCols; 

__global__ void FilterKernel (unsigned short *d_input_img, unsigned short *d_output_img, int d_iRows, int d_iCols)

{  
    unsigned short window[BLOCK_ROWS*BLOCK_COLS][KERNEL_DIMENSION*KERNEL_DIMENSION];

    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;   

    unsigned int tid = threadIdx.y*blockDim.y+threadIdx.x;

    if(x>d_iCols || y>d_iRows)
        return;

    window[tid][0]= (y==0||x==0) ? 0.0f : d_input_img[(y-1)*d_iCols+(x-1)];
    window[tid][1]= (y==0) ? 0.0f : d_input_img[(y-1)*d_iCols+x]; 
    window[tid][2]= (y==0||x==d_iCols-1) ? 0.0f : d_input_img[(y-1)*d_iCols+(x+1)];
    window[tid][3]= (x==0) ? 0.0f : d_input_img[y*d_iCols+(x-1)];
    window[tid][4]= d_input_img[y*d_iCols+x];
    window[tid][5]= (x==d_iCols-1) ? 0.0f : d_input_img[y*d_iCols+(x+1)];
    window[tid][6]= (y==d_iRows-1||x==0) ? 0.0f : d_input_img[(y+1)*d_iCols+(x-1)];
    window[tid][7]= (y==d_iRows-1) ? 0.0f : d_input_img[(y+1)*d_iCols+x];
    window[tid][8]= (y==d_iRows-1||x==d_iCols-1) ? 0.0f : d_input_img[(y+1)*d_iCols+(x+1)];

   __syncthreads();

    // Order elements 
    for (unsigned int j=0; j<9; ++j)
    {
        // Find position of minimum element
        int min=j;
        for (unsigned int l=j+1; l<9; ++l)
            if (window[tid][l] < window[tid][min])
                min=l;

        // Put found minimum element in its place
        const unsigned char temp=window[tid][j];
        window[tid][j]=window[tid][min];
        window[tid][min]=temp;

        __syncthreads();
    }

    d_output_img[y*d_iCols + x] = (window[tid][4]);

}

void take_input(const cv::Mat& input, const cv::Mat& output)
{

    unsigned short *device_input; 
    unsigned short *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_ROWS, BLOCK_COLS);  // 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( 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));

}

основная функция

#pragma once
#include<iostream>
#include<opencv2/core/core.hpp>
#include<opencv2/highgui/highgui.hpp>
#include<opencv2/imgproc/imgproc.hpp>
#include<opencv2/gpu/gpu.hpp>

#include <CUDA_wrapper.h>

using std::cout;
using std::endl;

int main()
{   

    //Read the image from harddisk, into a cv::Mat
    //IplImage *img=cvLoadImage("image.jpg");
    //cv::Mat input(img);
    cv::Mat input = cv::imread("C:/Users/OCT/Documents/Visual Studio 2008/Projects/MedianFilter/MedianFilter/pic1.bmp",CV_LOAD_IMAGE_GRAYSCALE);

    //IplImage* input = cvLoadImage("G:/Research/CUDA/Trials/OCTFilter/Debug/pic1.bmp");
    if(input.empty())
    {
        cout<<"Image Not Found"<<endl;
        getchar();
        return -1;
    }

    cv::Mat output(input.rows,input.cols,CV_8UC1);

    // store the different details of the input image like img_data, rows, cols in variables 
    int Rows = input.rows;
    int Cols = input.cols;
    unsigned char* Data = input.data;

    cout<<"image rows "<<Rows<<endl;
    cout<<"image cols "<<Cols<<endl;
    cout<<"\n"<<endl;
    cout<<"data "<<(int)Data<<endl;
    cv::waitKey(0);

    // call the device function to take the image as input
    take_input(input, output);

    cv::Mat dest; 

    medianBlur ( input, dest, 3 );

    //Show the input and output
    cv::imshow("Input",input);
    cv::imshow("Output",output);
    cv::imshow("Median blur",dest);

    //Wait for key press
    cv::waitKey();
}