CUDA：ループカーネルに埋め込まれています

Question

私はいくつかのコードをcudaカーネルに作りたいと思っています。見よ：

for (r = Y; r < Y + H; r+=2) 
    { 
     ch1RowSum = ch2RowSum = ch3RowSum = 0; 
     for (c = X; c < X + W; c+=2) 
     { 
      chan1Value = //some calc'd value 
          chan3Value = //some calc'd value 
      chan2Value = //some calc'd value 
      ch2RowSum += chan2Value; 
      ch3RowSum += chan3Value; 
      ch1RowSum += chan1Value; 
     } 
     ch1Mean += ch1RowSum/W; 
     ch2Mean += ch2RowSum/W; 
     ch3Mean += ch3RowSum/W; 
    } 

はこれは、一つの手段を計算するRowSumsと1を計算し、そしてどのように私は私のループインデックスはゼロから始まり、Nで終了いけないという事実を処理する必要があり、2つのカーネルに分割すべきか？

Answer 1

3つの値を計算するカーネルがあるとしましょう。構成内の各スレッドは、各（r、c）ペアの3つの値を計算します。

__global__ value_kernel(Y, H, X, W) 
{ 
    r = blockIdx.x + Y; 
    c = threadIdx.x + W; 

    chan1value = ... 
    chan2value = ... 
    chan3value = ... 
} 

上記のカーネルで合計（完全に並列）を計算することはできません。上記のように+ =を使用することはできません。そして、あなたは、各ブロック（行）に一つだけのスレッドがある場合は...このように、1つのカーネルに

__global__ both_kernel(Y, H, X, W) 
{ 
    r = blockIdx.x + Y; 
    c = threadIdx.x + W; 

    chan1value = ... 
    chan2value = ... 
    chan3value = ... 

    if(threadIdx.x == 0) 
    { 
     ch1RowSum = 0; 
     ch2RowSum = 0; 
     ch3RowSum = 0; 

     for(i=0; i<blockDim.x; i++) 
     { 
      ch1RowSum += chan1value; 
      ch2RowSum += chan2value; 
      ch3RowSum += chan3value; 
     } 

     ch1Mean = ch1RowSum/blockDim.x; 
     ch2Mean = ch2RowSum/blockDim.x; 
     ch3Mean = ch3RowSum/blockDim.x; 
    } 
} 

をそれをすべて合計しないと意味入れることができますが、それは最初の値のカーネルを使用するために、おそらくより良いですし、合計と手段の両方のための第2のカーネル...下のカーネルをさらに並列化することができます。別のカーネルを使用すると、準備が整ったらそれに集中することができます。

__global__ sum_kernel(Y,W) 
{ 
    r = blockIdx.x + Y; 

    ch1RowSum = 0; 
    ch2RowSum = 0; 
    ch3RowSum = 0; 

    for(i=0; i<W; i++) 
    { 
     ch1RowSum += chan1value; 
     ch2RowSum += chan2value; 
     ch3RowSum += chan3value; 
    } 

    ch1Mean = ch1RowSum/W; 
    ch2Mean = ch2RowSum/W; 
    ch3Mean = ch3RowSum/W; 
}