Sobel filter in cuda (can not show full image)

Question

У меня есть классическая проблема с выходом фильтра sobel с использованием CUDA.

это основной класс (main.cpp)

/*main class */ 
int main(int argc, char** argv) 
{ 
    IplImage* image_source = cvLoadImage("test.jpg",  
        CV_LOAD_IMAGE_GRAYSCALE); 
    IplImage* image_input = cvCreateImage(cvGetSize(image_source), 
         IPL_DEPTH_8U,image_source->nChannels); 
    IplImage* image_output = cvCreateImage(cvGetSize(image_source), 
         IPL_DEPTH_8U,image_source->nChannels); 

    /* Convert from IplImage tofloat */ 
    cvConvert(image_source,image_input); 

    unsigned char *h_out = (unsigned char*)image_output->imageData; 
    unsigned char *h_in = (unsigned char*)image_input->imageData; 

width  = image_input->width; 
height = image_input->height; 
widthStep = image_input->widthStep; 

sobel_parallel(h_in, h_out, width, height, widthStep); 
cvShowImage("CPU", image_output); 
cvReleaseImage(&image_output); 
waitKey(0); 
} 

И это файл CUDA (kernel_gpu.cu)

__global__ void kernel (unsigned char *d_in , unsigned char *d_out , int width , 
    int height, int widthStep) { 

int col = blockIdx . x * blockDim . x + threadIdx . x ; 
int row = blockIdx . y * blockDim . y + threadIdx . y ; 

int dx [3][3] = { -1 , 0 , 1 , 
       -2 , 0 , 2 , 
       -1 , 0 , 1}; 

int dy [3][3] = {1 ,2 ,1 , 
       0 ,0 ,0 , 
       -1 , -2 , -1}; 

int s; 
if(col < width && row < height) 
    { 
     int i = row; 
     int j = col; 
     // apply kernel in X direction 
     int sum_x=0; 
     for(int m=-1; m<=1; m++) 
      for(int n=-1; n<=1; n++) 
      { 
      s=d_in[(i+m)*widthStep+j+n]; // get the (i,j) pixel value 
      sum_x+=s*dx[m+1][n+1]; 
      } 
     // apply kernel in Y direction 
     int sum_y=0; 
     for(int m=-1; m<=1; m++) 
      for(int n=-1; n<=1; n++) 
      { 
      s=d_in[(i+m)*widthStep+j+n]; // get the (i,j) pixel value 
      sum_y+=s*dy[m+1][n+1]; 
      } 
     int sum=abs(sum_x)+abs(sum_y); 
     if (sum>255) 
      sum=255; 
     d_out[i*widthStep+j]=sum; // set the (i,j) pixel value 
    } 

} 
// Kernel Calling Function 

extern "C" void sobel_parallel(unsigned char* h_in, unsigned char* h_out, 
    int rows, int cols, int widthStep){ 

unsigned char* d_in; 
unsigned char* d_out; 
cudaMalloc((void**) &d_in, rows*cols); 
cudaMalloc((void**) &d_out, rows*cols); 

cudaMemcpy(d_in, h_in, rows*cols*sizeof(unsigned char), cudaMemcpyHostToDevice); 
dim3 block (16,16); 
dim3 grid ((rows * cols)/256.0); 
    kernel<<<grid,block>>>(d_in, d_out, rows, cols, widthStep); 

cudaMemcpy(h_out, d_out, rows*cols*sizeof(unsigned char), cudaMemcpyDeviceToHost); 
cudaFree(d_in); 
cudaFree(d_out); 
} 

Ошибка: результат изображение не появляется в их полностью, только часть изображения.

Почему результат (GPU) подобен этому? (Я попытался вычислить CPU с использованием той же функции и без проблем).

Answer 1

Вы создаете 1 мерную сетку, используя 2D-индексирование внутри ядра, которое будет охватывать только направление x, и будут отфильтрованы только верхние 16 строк изображения (потому что высота блока равна 16).

dim3 grid ((rows * cols)/256.0); //This is incorrect in current case 

Рассмотрите возможность создания 2-мерной сетки, чтобы она охватывала все строки изображения.

dim3 grid ((cols + 15)/16, (rows + 15)/16); 

Answer 2

Проверьте переменные width и widthStep, чтобы убедиться, что они на самом деле равны или нет, потому что в вашей функции sobel_parallel вы неявно предполагаете это (что может быть не так, поскольку ваши данные выровнены). Если это не так Кодекса

cudaMalloc((void**) &d_in, rows*cols); 

фактически выделять меньше памяти, чем это необходимо, и, следовательно, вы будете обрабатывать только часть изображения. Было бы лучше использовать

cudaMalloc((void**) &d_in, rows*widthStep); 

И, конечно, при необходимости отрегулируйте остальную часть кода.

Вы также призывают

void sobel_parallel(unsigned char* h_in, unsigned char* h_out, 
int rows, int cols, int widthStep) 

с

sobel_parallel(h_in, h_out, width, height, widthStep); 

, который обменивает строки с перевалы, и это снова обменялись, когда вы вызываете ядро. Это вызовет проблему, если вы используете приведенное выше предложение.