2012-04-27 24 views
1

現在、CUDAの配列シフト操作を行っていますが、私はGPUの操作を並列化する必要がある部分に固執しています。したがって、操作は基本的に配列内の要素を移動しています。CUDA配列要素のシフト操作

たとえば、M行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が表示されます

これは基本的に操作ですが、私の質問はどのように各行にスレッドを割り当てるのですか? またはすべての行を並列化して同じようにシフトしますクダの時間?また、私の 配列は、以下の式に

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

を使用して1次元配列に2次元配列から変換されたこれまでのところ私はこれを試してみた:

__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] : 
} 
} 

それの出力が正確に正しい(オフではありません私は間違って何をしているのか分かりません。

+0

あなたは1行につき-1以上を持っていますか?シフト後の右の値を埋めるためのルールは何ですか?あなたはいつも最後の値を複製しますか?列に-1が複数ある場合はどうなりますか? – harrism

+0

OPはすでに-1が行ごとに1回だけ発生することを明確にしました。そして、このスキームは、-1の後ろのすべてに移動するように思われますが、右の要素をそのまま残します。 –

答えて

1

これは、「述語和」をプリミティブとして使用するわずかに変更された「ストリームコンパクション」アルゴリズムで実行できるようです。詳細については、次のリンクを参照してください。 http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html

Hmm。 (コピー元のデータと-1を比較して)投票関数を使用して、ワープスレッドがコピーを実行するときにどのようにワープスレッドが宛先オフセットを選択するかを判断することに利点があることがわかります。