2016-07-31 27 views
-1

Ich habe eine Frage über Image Convolution in CUDA. Wenn ich es mit einer kleinen Maxtrix (16 * 16) teste, ist alles in Ordnung. Aber bei einer größeren Matrix ändert sich das Ergebnis immer, wenn ich laufe. Ich denke, Problem ist 2 für Schleifen in den Kernel.Wie kann ich Convolution Image in CUDA

__global__ void image_convolution_kernel(float *input, float *out, float *kernelConv, 
        int img_width, const int img_height, 
        const int kernel_width, const int kernel_height) 
{ 

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

    float sum = 0; 
    for (int j = 0; j < kernel_height; j++) 
    { 
     for (int i = 0; i < kernel_width; i++) 
     { 
      int dX = x + i - kernel_width/2; 
      int dY = y + j - kernel_height/2; 

      if (dX < 0) 
       dX = 0; 

      if (dX >= img_width) 
       dX = img_width - 1; 

      if (dY < 0) 
       dY = 0; 

      if (dY >= img_height) 
       dY = img_height - 1; 


      const int idMat = j * kernel_width + i; 
      const int idPixel = dY * img_width + dX; 
      sum += (float)input[idPixel] * kernelConv[idMat]; 
     } 
    } 

    const int idOut = y * img_width + x; 
    out[idOut] = abs(sum); 

} 

    void image_convolution(float * input,float* output, int img_height, int img_width) 
{ 
    int kernel_height = 3; 
    int kernel_width = 3; 
    float kernel[] ={ 0,-0.25,0, 
        -0.25,1,-0.25, 
         0,-0.25,0 
        }; 
    float * mask = new float[kernel_height*kernel_width]; 
    for (int i = 0; i < kernel_height*kernel_width; i++) 
    { 
     mask[i] = kernel[i]; 
    } 

    float * d_input, * d_output, * d_kernel; 
    cudaMalloc(&d_input, img_width*img_height*sizeof(float)); 
    cudaMalloc(&d_output, img_width*img_height*sizeof(float)); 
    cudaMalloc(&d_kernel, kernel_height*kernel_width*sizeof(float)); 

    cudaMemcpy(d_input, input, img_width*img_height*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_kernel, mask, kernel_height*kernel_width*sizeof(float), cudaMemcpyHostToDevice); 
    dim3 blocksize(16,16); 
    dim3 gridsize; 
    gridsize.x=(img_width+blocksize.x-1)/blocksize.x; 
    gridsize.y=(img_height+blocksize.y-1)/blocksize.y; 
    image_convolution_kernel<<<gridsize,blocksize>>>(d_input,d_output,d_kernel,img_width,img_height,kernel_width,kernel_height); 
    cudaMemcpy(output, d_output, img_width*img_height*sizeof(float), cudaMemcpyDeviceToHost); 

    for (int i=0; i < img_width*img_height; i++) 
    { 
     printf("%d, ",(int)output[i]); 
    } 
    printf("\n\n"); 
} 

Hier mein Ergebnis ist, teste ich es mit 24 * 24 Bild, ich laufe es zwei Mal, und ich schreibe auch einfache Funktion zum Vergleich der Ausgabe.

enter image description here

Und hier ist das Ergebnis, wenn ich die Ausgabe vergleichen, gibt es 32 verschiedene, bei Index 240, 241 .... enter image description here

+0

@RobertCrovella Sorry, ich wurde den Code aktualisiert und mit Bildern ausgegeben, bitte geben Sie mir einige Hinweise. – Jim

+0

Wenn Sie textbasierte Daten in eine Frage einfügen möchten, wird vorgeschlagen, dass Sie den tatsächlichen Text kopieren und einfügen (und ihn formatieren!), Anstatt Bilder des Texts in Ihre Frage einzufügen. –

Antwort

1

Sie einen ziemlich häufigen Fehler in Ihrem Programm gemacht hat. Beim Erstellen eines Raster von Fäden wie folgt aus:

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

Sie absichtlich zu schaffen (in der Regel) zusätzliche Fäden in jeder Dimension, um das Problem vollständig zu bedecken Raum (d.h. die Bildgröße). Daran ist nichts falsch.

Es bedeutet jedoch, dass wir extra Threads starten werden, die außerhalb der gültigen Bilddimension sind. Wir müssen sicherstellen, dass diese Threads nichts tun. Der übliche Ansatz besteht darin, dem Kernel eine Thread-Überprüfung hinzuzufügen, sodass Threads außerhalb der gültigen Image-Dimensionen nichts tun. Hier ist eine modifizierte Kernel und voll gearbeitet Beispiel zeigt, dass der Wandel:

$ cat t1219.cu 
#include <iostream> 
#include <cstdlib> 

const int iw = 1025; 
const int ih = 1025; 
const int rng = 10; 

__global__ void image_convolution_kernel(float *input, float *out, float *kernelConv, 
        int img_width, const int img_height, 
        const int kernel_width, const int kernel_height) 
{ 

    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    if ((x < img_width) && (y < img_height)){ // thread check 
     float sum = 0; 
     for (int j = 0; j < kernel_height; j++) 
     { 
     for (int i = 0; i < kernel_width; i++) 
     { 
      int dX = x + i - kernel_width/2; 
      int dY = y + j - kernel_height/2; 

      if (dX < 0) 
       dX = 0; 

      if (dX >= img_width) 
       dX = img_width - 1; 

      if (dY < 0) 
       dY = 0; 

      if (dY >= img_height) 
       dY = img_height - 1; 


      const int idMat = j * kernel_width + i; 
      const int idPixel = dY * img_width + dX; 
      sum += (float)input[idPixel] * kernelConv[idMat]; 
     } 
     } 

     const int idOut = y * img_width + x; 
     out[idOut] = abs(sum); 
    } 

} 

    void image_convolution(float * input,float* output, int img_height, int img_width) 
{ 
    int kernel_height = 3; 
    int kernel_width = 3; 
    float kernel[] ={ 0,-0.25,0, 
        -0.25,1,-0.25, 
         0,-0.25,0 
        }; 
    float * mask = new float[kernel_height*kernel_width]; 
    for (int i = 0; i < kernel_height*kernel_width; i++) 
    { 
     mask[i] = kernel[i]; 
    } 

    float * d_input, * d_output, * d_kernel; 
    cudaMalloc(&d_input, img_width*img_height*sizeof(float)); 
    cudaMalloc(&d_output, img_width*img_height*sizeof(float)); 
    cudaMalloc(&d_kernel, kernel_height*kernel_width*sizeof(float)); 

    cudaMemcpy(d_input, input, img_width*img_height*sizeof(float), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_kernel, mask, kernel_height*kernel_width*sizeof(float), cudaMemcpyHostToDevice); 
    dim3 blocksize(16,16); 
    dim3 gridsize; 
    gridsize.x=(img_width+blocksize.x-1)/blocksize.x; 
    gridsize.y=(img_height+blocksize.y-1)/blocksize.y; 
    image_convolution_kernel<<<gridsize,blocksize>>>(d_input,d_output,d_kernel,img_width,img_height,kernel_width,kernel_height); 
    cudaMemcpy(output, d_output, img_width*img_height*sizeof(float), cudaMemcpyDeviceToHost); 
} 

int main(){ 

    float *in, *out; 
    int is = ih*iw; 
    in = new float[is]; 
    out = new float[is]; 
    for (int i = 0; i < is; i++) {in[i] = rand()%rng; out[i] = -1;} 
    image_convolution(in,out, ih, iw); 
    for (int iy = 1; iy < ih-1; iy++) 
    for (int ix = 1; ix < iw-1; ix++){ 
     float temp = abs(-0.25 * (in[iy*iw + ix -1] + in[iy*iw + ix +1] + in[(iy-1)*iw + ix] + in[(iy+1)*iw + ix]) + in[iy*iw+ix]); 
     if (out[iy*iw+ix] != temp) {std::cout << "mismatch x: " << ix << " y: " << iy << " was: " << out[iy*iw+ix] << " should be: " << temp << std::endl; return 1;}} 
    return 0; 
} 
$ nvcc -o t1219 t1219.cu 
$ cuda-memcheck ./t1219 
========= CUDA-MEMCHECK 
========= ERROR SUMMARY: 0 errors 
$ 

Für die Bildabmessungen, die genaue Vielfaches der Blockgröße sind (16,16) wurde dieses Problem (die für meine vorherigen Testfall wahr war) nicht auftauchen - der Code wird korrekt funktionieren. Für alle anderen Testfälle benötigen wir eine solche Thread-Prüfung.

+0

Danke für deine Hilfe, ich habe das Problem gelöst :) – Jim