3

CUDA を使用して特定の画像処理アルゴリズムを実装していますが、スレッド同期の問題全体についていくつか質問があります。

当面の問題は次のように説明できます。

W*H のサイズの画像があります。画像の各ピクセルに対して、9 つの同一データの並列プロセスを実行する必要があり、各プロセスは結果として値の配列を返します (配列はアルゴリズム全体で同じ長さです。N は約 20 または 30 です)。 ）。ピクセルごとに、これらの 9 つのプロセスは、計算を終了した後、結果を最終的な配列 (ピクセルごとに 1 つの配列) に蓄積します。

これを並列化するために、次の構造を設計しました: (10,10,9) の次元のブロックを生成します。これは、各スレッド ブロックが 10*10 サイズのサブ イメージを処理し、各スレッドが 9 つのうちの 1 つを処理することを意味します。単一のピクセルに対して同じプロセス。この場合、グリッドの寸法は (W/10,H/10,1) になります。スレッド ブロックの場合、長さ 100*N の共有メモリ配列を割り当て、各スレッドは現在のピクセルの座標に従って適切な共有メモリ位置に書き込みます。したがって、ここでは、atomicAdd と __synchthreads() との同期が必要です。

ここでの問題は、ピクセルの値がゼロの場合、それを処理する必要がまったくないため、そのようなピクセルを終了したいということです。そうしないと、画像の大部分がゼロ (背景)。ということで、以下のように書こうと思いました。

//X and Y are the coordinates of the current pixel in the input image.
//threadIdx.z gives the index of the process among the 9 for the current pixel. 

int X=blockIdx.x * blockDim.x + threadIdx.x;
int Y=blockIdx.y * blockDim.y + threadIdx.y;
int numOfProcessForTheCurrPixel=threadIdx.z;
int linearIndexOfPixelInBlock=threadIdx.y * blockDim.x + threadIdx.x;

unsigned short pixelValue=tex2D(image,X,Y);
//Here, threads processing zero-pixels will exit immediately.
if(pixelValue==0)
 return;

float resultArray[22];
//Fill the result array according to our algorithm, mostly irrelevant stuff.
ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

for(int i=0;i<22;i++)
    atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

 __syncthreads(); 
 //Then copy from the shared to the global memory and etc.

この状況で私が心配しているのは、プログラミングガイドが言っていることです:

__syncthreads() は条件付きコードで使用できますが、条件がスレッド ブロック全体で同じように評価される場合に限ります。そうしないと、コードの実行がハングしたり、意図しない副作用が発生したりする可能性があります。

したがって、私の場合、10*10 スレッド ブロック内のピクセルの一部がゼロであるかどうかにかかわらず、ゼロ ピクセルに属するスレッドは最初にすぐに終了し、他のスレッドは処理を続行します。この場合、同期はどうなりますか? プログラミング ガイドにあるように、同期は正常に動作しますか? それとも未定義の動作を生成しますか? ゼロ ピクセル スレッドでガベージ データを処理してビジー状態に保つことを考えましたが、完全にゼロで構成されるブロックがある場合 (非常に頻繁に発生する)、処理時間が不必要に長くなります。この場合、何ができるでしょうか？

4

1 に答える 1

1

デッドロックの発生を回避するには、すべてのスレッドが無条件に_synchthreads（）にヒットする必要があります。この例では、returnを、関数の大部分を飛び越えて、ゼロピクセルの場合の_syncthreads（）に直接向かうifステートメントに置き換えることでこれを行うことができます。

unsigned short pixelValue=tex2D(image,X,Y);
//If there's nothing to compute, jump over all the computation stuff
if(pixelValue!=0)
{

    float resultArray[22];
    //Fill the result array according to our algorithm, mostly irrelevant stuff.
    ProcessPixel(resultArray,X,Y,numOfProcessForTheCurrPixel);

    for(int i=0;i<22;i++)
        atomicAdd(&__sharedMemoryArray[22*linearIndexOfPixelInBlock + i],resultArray[i]);

}

__syncthreads(); 

if (pixelValue != 0)
{
    //Then copy from the shared to the global memory and etc. 
}
于 2012-07-25T17:14:30.743 に答える