image - 使用 CUDA 和 OpenCV 在灰度 jpg 图像上应用 Sobel 边缘检测

Question

这个问题之前已经有人问过，但是提问者没有提供足够的信息而没有得到回答，我对这个程序很好奇。

Original Question Link

我正在尝试使用 opencv 和 cuda 库进行 sobel 边缘检测，X方向的索贝尔核是

-1 0 1   
-2 0 2  
-1 0 1   

我的项目中有3个文件

main.cpp
CudaKernel.cu
CudaKernel.h

main.cpp

#include <stdlib.h>
#include <iostream>
#include <string.h>
#include <Windows.h>
#include <opencv2\core\core.hpp>
#include <opencv2\highgui\highgui.hpp>
#include <opencv2\gpu\gpu.hpp>
#include <cuda_runtime.h>
#include <cuda_gl_interop.h>
#include "CudaKernel.h"

using namespace cv;
using namespace std;


int main(int argc, char** argv) 
{
    IplImage* image;

    try
    {
        image = cvLoadImage("4555472_460s.jpg", CV_LOAD_IMAGE_GRAYSCALE);
        gpu::DeviceInfo info = gpu::getDevice();
        cout << info.name() << endl;
        cout << "Stream Processor : "<< info.multiProcessorCount() << endl;
        cout << "Total Graphic Memory :" << info.totalMemory()/1048576 << " MB" << endl; 
    }
    catch (const cv::Exception* ex)
    {
        cout << "Error: " << ex->what() << endl;
    }
    if(!image )
        {
             cout << "Could not open or find the image" << std::endl ;
             return -1;
        }


    IplImage* image2=cvCreateImage(cvGetSize(image),IPL_DEPTH_32F,image->nChannels);
    IplImage* image3=cvCreateImage(cvGetSize(image),IPL_DEPTH_32F,image->nChannels);

    unsigned char * pseudo_input=(unsigned char *)image->imageData;
    float *output=(float*)image2->imageData;
    float *input=(float*)image3->imageData;
    int s=image->widthStep/sizeof(float);
        for(int w=0;w<=(image->height);w++)
            for(int h=0;h<(image->width*image->nChannels);h++)
            {
                input[w*s+h]= pseudo_input[w*s+h];
            }


    Pixel *fagget  = (unsigned char*) image->imageData;
    kernelcall(input, output, image->width,image->height, image->widthStep);

//  cv::namedWindow( "Display window", CV_WINDOW_AUTOSIZE );// Create a window for display.
    cvShowImage( "Original Image", image ); // Show our image inside it.
    cvShowImage("Sobeled Image", image2);
    waitKey(0); // Wait for a keystroke in the window
    return 0;

}

CudaKernel.cu

#include<cuda.h>
#include<iostream>
#include "CudaKernel.h"
using namespace std;
#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
#define CudaCheckError()    __cudaCheckError( __FILE__, __LINE__ )
#define checkCudaErrors(err) __checkCudaErrors (err, __FILE__, __LINE__)


texture <float,2,cudaReadModeElementType> tex1;
texture<unsigned char, 2> tex;
static cudaArray *array = NULL;
static cudaArray *cuArray = NULL;


//Kernel for x direction sobel
__global__ void implement_x_sobel(float* garbage,float* output,int width,int height,int widthStep)
{
    int x=blockIdx.x*blockDim.x+threadIdx.x;
    int y=blockIdx.y*blockDim.y+threadIdx.y;

    float output_value=((0*tex2D(tex1,x,y))+(2*tex2D(tex1,x+1,y))+(-2*tex2D(tex1,x-  1,y))+(0*tex2D(tex1,x,y+1))+(1*tex2D(tex1,x+1,y+1))+(-1*tex2D(tex1,x-1,y+1))+  (1*tex2D(tex1,x+1,y-1))+(0*tex2D(tex1,x,y-1))+(-1*tex2D(tex1,x-1,y-1)));
    output[y*widthStep+x]=output_value;
}


inline void __checkCudaErrors( cudaError err, const char *file, const int line )
{
    if( cudaSuccess != err) {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",
                file, line, (int)err, cudaGetErrorString( err ) );
        exit(-1);
    }
}   

//Host Code
 inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
if ( cudaSuccess != err )
{
    printf("cudaSafeCall() failed at %s:%i : %s\n",
             file, line, cudaGetErrorString( err ) );
    exit( -1 );
}    
#endif

return;
}
inline void __cudaCheckError( const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
cudaError err = cudaGetLastError();
if ( cudaSuccess != err )
{
    printf("cudaCheckError() failed at %s:%i : %s\n",
             file, line, cudaGetErrorString( err ) );
   exit( -1 );
}
#endif

return;
}

void kernelcall(float* input,float* output,int width,int height,int widthStep){
    //cudaChannelFormatDesc channelDesc=cudaCreateChannelDesc(32,32,0,0,cudaChannelFormatKindFloat);
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
    //cudaArray *cuArray;
    CudaSafeCall(cudaMallocArray(&cuArray,&channelDesc,width,height));
    cudaMemcpyToArray(cuArray,0,0,input,widthStep*height,cudaMemcpyHostToDevice);

    tex1.addressMode[0]=cudaAddressModeClamp;
    tex1.addressMode[1]=cudaAddressModeClamp;
    tex1.filterMode=cudaFilterModeLinear;
    cudaBindTextureToArray(tex1,cuArray,channelDesc);
    tex1.normalized=false;
    float * D_output_x;
    float * garbage=NULL;
    CudaSafeCall(cudaMalloc(&D_output_x,widthStep*height)); 
    dim3 blocksize(16,16);
    dim3 gridsize;
    gridsize.x=(width+blocksize.x-1)/blocksize.x;
    gridsize.y=(height+blocksize.y-1)/blocksize.y;

    implement_x_sobel<<<gridsize,blocksize>>>(garbage,D_output_x,width,height,widthStep/sizeof(float));
    cudaThreadSynchronize();
    CudaCheckError();
    CudaSafeCall(cudaMemcpy(output,D_output_x,height*widthStep,cudaMemcpyDeviceToHost));
    cudaFree(D_output_x);
    cudaFree(garbage);
    cudaFreeArray(cuArray);
}

结果真的很乱，跟原图完全不一样

结果:

我将一些代码更改为

float *pseudo_input=(float *)image->imageData;
float *output=(float*)image2->imageData;
float *input=(float*)image3->imageData;
float *inputnormalized=(float *)image4->imageData;

int s=image->widthStep/sizeof(float);
for(int w=0;w<=(image->height);w++)
    for(int h=0;h<(image->width*image->nChannels);h++)
    {
        input[w*s+h]= pseudo_input[w*s+h];
    }


kernelcall(input, output, image->width,image->height, image->widthStep);

cvNormalize(input,inputnormalized,0,255,NORM_MINMAX, CV_8UC1);

cvShowImage( "Original Image", image ); // Show our image inside it.
cvShowImage("Sobeled Image", image2);

但现在我收到一个未处理的异常错误。

Answer 1

OpenCV 规则编号 1:

Never access the image data directly through the underlying data pointer unless absolutely necessary, e.g copying data to GPU. Reference (Me :p)

错误/建议:

1. 不是通过循环遍历图像数据来转换图像指针，使用 cvConvert 改变图像数据类型。循环很很容易出错。

2. 当调用名为kernelcall 的函数时，您正在传递float 图像的数据指针，但传递 widthStep原始 8 位图像。这是错误结果的主要原因，因为它将导致内核内部的索引不正确。

3. 在 2 个倾斜指针之间执行内存复制时不同的 widthSteps，始终使用可用的 2D 内存复制功能在 CUDA 运行时，例如cudaMemcpy2D、cudaMemcpy2DToArray 等。在您的情况下，cuArray 内部具有未知的宽度步长，并且输入 IplImage 具有不同的widthStep 比 cuArray .

4. 避免不必要的 header 、分配和标识符声明。

5. 在 CUDA 内核中添加边界检查，以便只有那些线程执行落在图像中的内存读/写。它可能会导致一点分歧，但它比无效的内存读/写要好。

修改后的代码(已测试):

main.cpp

#include <iostream>
#include <opencv2/opencv.hpp>
#include "CudaKernel.h"

using namespace cv;
using namespace std;

int main(int argc, char** argv) 
{
    IplImage* image;

    image = cvLoadImage("4555472_460s.jpg", CV_LOAD_IMAGE_GRAYSCALE);

    if(!image )
    {
        cout << "Could not open or find the image" << std::endl;
        return -1;
    }


    IplImage* image2 = cvCreateImage(cvGetSize(image),IPL_DEPTH_32F,image->nChannels);
    IplImage* image3 = cvCreateImage(cvGetSize(image),IPL_DEPTH_32F,image->nChannels);

    //Convert the input image to float
    cvConvert(image,image3);

    float *output = (float*)image2->imageData;
    float *input =  (float*)image3->imageData;

    kernelcall(input, output, image->width,image->height, image3->widthStep);

    //Normalize the output values from 0.0 to 1.0
    cvScale(image2,image2,1.0/255.0);

    cvShowImage("Original Image", image );
    cvShowImage("Sobeled Image", image2);
    cvWaitKey(0);
    return 0;
}

CudaKernel.cu

#include<cuda.h>
#include<iostream>
#include "CudaKernel.h"

using namespace std;

#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
#define CudaCheckError()    __cudaCheckError( __FILE__, __LINE__ )
#define checkCudaErrors(err) __checkCudaErrors (err, __FILE__, __LINE__)


texture <float,2,cudaReadModeElementType> tex1;

static cudaArray *cuArray = NULL;

//Kernel for x direction sobel
__global__ void implement_x_sobel(float* output,int width,int height,int widthStep)
{
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    //Make sure that thread is inside image bounds
    if(x<width && y<height)
    {
        float output_value = (-1*tex2D(tex1,x-1,y-1)) + (0*tex2D(tex1,x,y-1)) + (1*tex2D(tex1,x+1,y-1))
                           + (-2*tex2D(tex1,x-1,y))   + (0*tex2D(tex1,x,y))   + (2*tex2D(tex1,x+1,y))
                           + (-1*tex2D(tex1,x-1,y+1)) + (0*tex2D(tex1,x,y+1)) + (1*tex2D(tex1,x+1,y+1));

        output[y*widthStep+x]=output_value;
    }

}


inline void __checkCudaErrors( cudaError err, const char *file, const int line )
{
    if( cudaSuccess != err) {
        fprintf(stderr, "%s(%i) : CUDA Runtime API error %d: %s.\n",
            file, line, (int)err, cudaGetErrorString( err ) );
        exit(-1);
    }
}   

//Host Code
inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
    if ( cudaSuccess != err )
    {
        printf("cudaSafeCall() failed at %s:%i : %s\n",
            file, line, cudaGetErrorString( err ) );
        exit( -1 );
    }    
#endif

    return;
}
inline void __cudaCheckError( const char *file, const int line )
{
#ifdef CUDA_ERROR_CHECK
    cudaError err = cudaGetLastError();
    if ( cudaSuccess != err )
    {
        printf("cudaCheckError() failed at %s:%i : %s\n",
            file, line, cudaGetErrorString( err ) );
        exit( -1 );
    }
#endif

    return;
}

void kernelcall(float* input,float* output,int width,int height,int widthStep)
{
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();

    CudaSafeCall(cudaMallocArray(&cuArray,&channelDesc,width,height));

    //Never use 1D memory copy if host and device pointers have different widthStep.
    // You don't know the width step of CUDA array, so its better to use cudaMemcpy2D...
    cudaMemcpy2DToArray(cuArray,0,0,input,widthStep,width * sizeof(float),height,cudaMemcpyHostToDevice);

    cudaBindTextureToArray(tex1,cuArray,channelDesc);

    float * D_output_x;
    CudaSafeCall(cudaMalloc(&D_output_x,widthStep*height)); 

    dim3 blocksize(16,16);
    dim3 gridsize;
    gridsize.x=(width+blocksize.x-1)/blocksize.x;
    gridsize.y=(height+blocksize.y-1)/blocksize.y;

    implement_x_sobel<<<gridsize,blocksize>>>(D_output_x,width,height,widthStep/sizeof(float));

    cudaThreadSynchronize();
    CudaCheckError();

    //Don't forget to unbind the texture
    cudaUnbindTexture(tex1);

    CudaSafeCall(cudaMemcpy(output,D_output_x,height*widthStep,cudaMemcpyDeviceToHost));

    cudaFree(D_output_x);
    cudaFreeArray(cuArray);
}