CUDA Cプログラムのifとelse条件の並列化

Question

Cで画像をぼかすフィルタを作成しましたが正常に動作しました.CUDA Cを使用して処理を高速化するためにGPUで実行しようとしています。このプログラムには、Cコードのバージョンである以下のようないくつかのifおよびelse条件があります。 関数への入力は、入力画像、出力画像、および列のサイズです。

void convolve_young1D(double * in, double * out, int datasize) { 
     int i, j; 

    /* Compute first 3 output elements */ 
    out[0] = B*in[0]; 
    out[1] = B*in[1] + bf[2]*out[0]; 
    out[2] = B*in[2] + (bf[1]*out[0]+bf[2]*out[1]); 

    /* Recursive computation of output in forward direction using filter parameters bf and B */ 
    for (i=3; i<datasize; i++) { 
     out[i] = B*in[i]; 
     for (j=0; j<3; j++) { 
      out[i] += bf[j]*out[i-(3-j)]; 
     } 
    } 

} 
    //Calling function below 
void convolve_young2D(int rows, int columns, int sigma, double ** ip_padded) { 

     /** \brief Filter radius */ 
     w = 3*sigma; 
     /** \brief Filter parameter q */ 
     double q; 
     if (sigma < 2.5) 
      q = 3.97156 - 4.14554*sqrt(1-0.26891*sigma); 
     else 
      q = 0.98711*sigma - 0.9633; 

     /** \brief Filter parameters b0, b1, b2, b3 */ 
     double b0 = 1.57825 + 2.44413*q + 1.4281*q*q + 0.422205*q*q*q; 
     double b1 = 2.44413*q + 2.85619*q*q + 1.26661*q*q*q; 
     double b2 = -(1.4281*q*q + 1.26661*q*q*q); 
     double b3 = 0.422205*q*q*q; 

     /** \brief Filter parameters bf, bb, B */ 
     bf[0] = b3/b0; bf[1] = b2/b0; bf[2] = b1/b0; 
     bb[0] = b1/b0; bb[1] = b2/b0; bb[2] = b3/b0; 
     B = 1 - (b1+b2+b3)/b0; 

     int i,j; 

     /* Convolve each row with 1D Gaussian filter */ 
     double *out_t = calloc(columns+(2*w),sizeof(double)); 
     for (i=0; i<rows+2*w; i++) { 
      convolve_young1D(ip_padded[i], out_t, columns+2*w); 
     } 
    free(out_t); 

は、CUDA C内のブロックとスレッドと同じアプローチを試みたが、私は出力としてゼロを取得してきたとも、入力値がゼロに変更するように見える成功しなかった私が間違っつもりです場所がわからないようにしてください助けてください。私はCUDA Cのプログラミングにはかなり新しいです。ここに私の試したバージョンのCUDAカーネルがあります。

__global__ void convolve_young2D(float *in, float *out,int rows,int columns, int j,float B,float bf[3],int w) { 

int k; 

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

if((x>0) && (x<(rows+2*w))) 
{ 
//printf("%d \t",x); 

if(j ==0) 
{ 
// Compute first output elements 
out[x*columns] = B*in[x*columns]; 
} 

else if(j==1) 
{ 
    out[x*columns +1 ] = B*in[x*columns +1] + bf[2]*out[x*columns]; 
} 
else if (j== 2) 
{ 
    out[2] = B*in[x*columns +2] + (bf[1]*out[x*columns]+bf[2]*out[x*columns+1]); 
} 
else{ 
// Recursive computation of output in forward direction using filter parameters bf and B 
     out[x*columns+j] = B*in[x*columns+j]; 
     for (k=0; k<3; k++) { 
      out[x*columns + j] += bf[k]*out[(x*columns+j)-(3-k)]; 
     } 
    } 
} 

} 

//Calling function below 
void convolve_young2D(int rows, int columns, int sigma, const float * const ip_padded, float * const op_padded) { 
float bf[3], bb[3]; 
float B; 
int w; 

    /** \brief Filter radius */ 

    w = 3*sigma; 
    /** \brief Filter parameter q */ 
    float q; 
    if (sigma < 2.5) 
     q = 3.97156 - 4.14554*sqrt(1-0.26891*sigma); 
    else 
     q = 0.98711*sigma - 0.9633; 
    /** \brief Filter parameters b0, b1, b2, b3 */ 
    float b0 = 1.57825 + 2.44413*q + 1.4281*q*q + 0.422205*q*q*q; 
    float b1 = 2.44413*q + 2.85619*q*q + 1.26661*q*q*q; 

    float b2 = -(1.4281*q*q + 1.26661*q*q*q); 
    float b3 = 0.422205*q*q*q; 
/** \brief Filter parameters bf, bb, B */ 
    bf[0] = b3/b0; bf[1] = b2/b0; bf[2] = b1/b0; 
    bb[0] = b1/b0; bb[1] = b2/b0; bb[2] = b3/b0; 
    B = 1 - (b1+b2+b3)/b0;  
    int p; 
const int inputBytes = (rows+2*w) * (columns+2*w) * sizeof(float); 
float *d_input, *d_output; // arrays in the GPU´s global memory 
cudaMalloc(&d_input, inputBytes); 
cudaMemcpy(d_input, ip_padded, inputBytes, cudaMemcpyHostToDevice); 
cudaMalloc(&d_output,inputBytes); 
for (p = 0; p<columns+2*w; p++){ 
convolve_young<<<4,500>>>(d_input,d_output,rows,columns,p,B,bf,w); 
} 
cudaMemcpy(op_padded, d_input, inputBytes, cudaMemcpyDeviceToHost); 
cudaFree(d_input); 

Answer 1

最初の問題は、あなたがconvolve_young<<<4,500>>>(d_input,d_output,rows,columns,p,B,bf,w);を呼んでいますが、2D convolve_youngという名前のカーネルを定義していることです。

別の可能性のある問題は、それはあなたが行う畳み込みを行うことです。

for (i=0; i<rows+2*w; i++) { 
    convolve_young1D(ip_padded[i], out_t, columns+2*w); 
} 

まず：あなたは、CPUのアルゴリズムに比べ行の代わりに列をループしている

ここ

for (p = 0; p<columns+2*w; p++){ 
    convolve_young<<<4,500>>>(d_input,d_output,rows,columns,p,B,bf,w); 
} 

あなたのCPUアルゴリズムの直接のポートを行い、その時に1行を計算してから、それを修正して画像全体を転送しなければならない。