小贝子编程

CUDA中值滤波器的实现不能产生期望的结果



我一直在尝试实现Wiki文章中提供的中值过滤器算法:http://en.wikipedia.org/wiki/Median_filter#2D_median_filter_pseudo_code

就我所知，我知道我所执行的是正确的。然而，当我查看结果时，我似乎无法获得与OpenCV中median blur函数产生的输出相似的输出。目前，我不关心通过使用共享内存或纹理内存来加速我的代码。我只是想先把事情办好。我的输入图像的大小是1024 x 256像素。

我做错了什么?是否有线程泄漏在我的代码?我知道我应该使用共享内存来防止全局读取，因为目前，我正在从全局内存中读取数据很多。

http://snag.gy/OkXzP.jpg——第一张图像是输入，第二张图像是我的算法结果，第三张是openCV medianblur函数结果。理想情况下，我希望我的算法输出与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 %dn", 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));
}

main函数

#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();
}

我相信在你的"内核实现"文件中有各种各样的错误和不必要的复杂性。

你可能会有更好的运气:

$ 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 %dn", 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: %dn", i, (int)input_im.data[i], (int)output_im.data[i]); return 1;}
  printf("Successn");
  return 0;
}

$ nvcc -o t376 t376.cu
$ ./t376
Success
$

注释:

  1. 我已经测试了这个(不使用OpenCV)对于我在代码中注入的简单情况。
    2. 你对window的使用太复杂了。注意，您使用它的方式是，每个线程都将实例化自己的本地window副本，独立于其他线程，并且对其他线程不可见。(也许您打算在这里使用共享内存?但是我离题了。)你的分类程序被打破了。我把它修改成我认为可以工作的版本。
  2. 将整个
    3. 中的像素数据类型替换为unsigned char
  3. x和y是混乱的，所以我把它们改为row和col，似乎不那么混乱。
  4. 稍微改进了内核错误检查
    5. 有很多方法可以优化这一点。然而，你明智的目标是首先让某些东西正常工作。所以我不会花太多时间在优化上，除了指出用于重用window数据的共享内存的两个一般区域，以及改进的排序例程。
  5. 你需要为openCV适当地修改这个
  6. 请注意，如果你修改它为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);

传递给imread的1参数指定RGB图像加载。参考imread文档:

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，它肯定不会工作。

在CUDA 6.0中，NPP库现在包括对所有数据类型和像素格式的中值过滤器的实现。如果你想要一个功能性的中值过滤器例程，你可以调用它。如果您需要帮助调试内核，请参阅所有其他答案…

相关内容

  • 没有找到相关文章

最新更新


热门标签：

javascript python java c# php android html jquery c++ css ios sql mysql arrays asp.net json python-3.x ruby-on-rails .net sql-server django objective-c excel regex ruby linux ajax iphone xml vba spring asp.net-mvc database wordpress string postgresql wpf windows xcode bash git oracle list vb.net multithreading eclipse algorithm macos powershell visual-studio image forms numpy scala function api selenium