2

現在、CUDAで配列シフト操作を実行していますが、GPUで操作を並列化する必要がある部分で立ち往生しています(すでにCPUで実行しています)。したがって、操作は基本的に配列内の要素をシフトすることです。

したがって、たとえば、M x N の行列がある場合、各行で -1 が表示された場合、-1 をその隣の要素に置き換えます。行の終わりに到達するまで、これを繰り返します。すべての列に対して並行してそれを行います。

簡単な例を次に示します。

 3  4  1 -1  5  6  7  8
-1  4  5  2  1  2  5  2
 2  4  5  1  2  3  4 -1

その行列の場合、結果の行列は次のようになります。

 3  4  1  5  6  7  8  8
 4  5  2  1  2  5  2  2
 2  4  5  1  2  3  4 -1

PS。最後の要素は、置き換えるものが何もない境界に到達するため、同じままです。また、各行に -1 が 1 つだけ表示されます。

だから、それは基本的に操作ですが、私の質問はどのように各行にスレッドを割り当てるか..すべての行を並列化し、cudaで同時にこのシフトを行うのですか? また、私の配列は、式を使用して2次元配列から1次元配列に変換されます

array1d[i+width*j]  =  array2d[i][j];

これまでのところ、私はこれを試しました:

__global__ void gpu_shiftArray(int *Arr, int *location, int width, int height)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;

int index = i+width*j;

//shift when I see -1
if(Arr[index] == -1)
{
    Arr[index] = (index % height) ? Arr[index+1] : 
    }
    //location stores the index of -1, so anything after the -1 will be shifted too
if((location[i]+width*j) <= index)
{
    Arr[index] = (index % height) ? Arr[index+1] : 
}
}

その出力は正確ではありません (5 ~ 10 の値がオフになっています) が、その理由が正確にはわかりませんし、何が間違っているのかもわかりません。

4

1 に答える 1

1

これは、「述語の合計」をプリミティブとして使用する、わずかに変更された「ストリーム圧縮」アルゴリズムで実行できるようです。詳細については、次のリンクを参照してください: Parallel Prefix Sum (Scan) with CUDA

うーん。バロット関数 (ソース データを -1 と比較する) といくつかのビット演算を使用して、ワープ スレッドがコピーを実行するときに宛先オフセットを選択する方法 (および選択するかどうか) を決定することにも利点があることがわかります。

于 2012-04-27T02:49:20.963 に答える