CUDA中x方向的Sobel滤波器

Question

我正在尝试在每个像素的 x 方向上的灰度图像上应用 sobel 过滤器并显示结果。 X方向sobel滤波器为：-

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

我没有得到所需的结果。有人可以指出我的错误吗？我正在尝试使用纹理，但我不确定我是否正确使用了它：

#include <cuda.h>
#include<iostream>
using namespace std;
#define CudaSafeCall( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
#define CudaCheckError()    __cudaCheckError( __FILE__, __LINE__ )
texture <float,2,cudaReadModeElementType> tex1;
//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;
}
//Kernel for y direction sobel
//__global__ void implement_y_sobel(float* input,float* output,int width,int height,int widthStep)
//{

//}
//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 sobel(float* input,float* output,int width,int height,int widthStep)
{
cudaChannelFormatDesc     channelDesc=cudaCreateChannelDesc(32,32,0,0,cudaChannelFormatKindFloat);
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;
tex1.normalized=false;
cudaBindTextureToArray(tex1,cuArray,channelDesc);
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;

//kernel call
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);
}

我的主文件：-

#include<iostream>
#include <stdio.h>
#include <stdlib.h>
#include<opencv/highgui.h>
#include<opencv/cv.h>
#include"header.h"
using namespace std;
void main()
{
IplImage* img1=cvLoadImage("C://test.jpg",CV_LOAD_IMAGE_GRAYSCALE); 
if( !img1) {
           printf("ERROR: couldnt load file!\n");
           }
IplImage* img2=cvCreateImage(cvGetSize(img1),IPL_DEPTH_32F,img1->nChannels);
IplImage* img3=cvCreateImage(cvGetSize(img1),IPL_DEPTH_32F,img1->nChannels);
unsigned char * pseudo_input=(unsigned char *)img1->imageData;
float * output=(float*)img2->imageData;
float *input=(float*)img3->imageData;
int s=img1->widthStep/sizeof(float);
for(int w=0;w<=(img1->height);w++)
    for(int h=0;h<(img1->width*img1->nChannels);h++)
    {
        input[w*s+h]= pseudo_input[w*s+h];
            }

sobel(input,output,img1->width,img1->height,img1->widthStep);
cvShowImage("Original Image",img1);
cvShowImage("Sobeled Image",img2);
cvWaitKey(0);
    }}

Answer 1

cudaCreateChannelDesc 期望 x、y、z 和 w 分量的位数作为前 4 个参数。 float 纹理应该是 32。

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 32, 0, 0, cudaChannelFormatKindFloat);

Answer 2

如果没有更多信息，很难诊断问题。如果您没有得到有意义的输出（例如纹理读取全 0），则意味着您的纹理设置或绑定存在问题。

如果你偏离了一点点，那可能是因为你需要将坐标偏移 0.5f，并且在你这样做的时候，在将你的整数显式转换为浮点数时要更加小心。如果您在调用 tex2D() 之前声明并分配浮点值变量，则代码不会运行得更慢。