0

グリッド内の制御スレッド ブロックについて質問があります。

私のソースは、画像の再帰的な仕事です。しかし、処理中、多​​くのブロックが約 8 回終了条件を満たしました。実行が 16 回以上ループする原因となったブロックはごくわずかでした。そこで、実行終了条件を満たしたブロックをスキップしたい。

可能です？

__global__ main(){
/* previous */
int *blockMap;
cudaMalloc((void**)&blockMap, sizeof(int) * nXBlockNum * nYBlockNum);
cudaMemset((void**)&blockMap, 0, sizeof(int) * nXBlockNum * nYBlockNum);

kernel<<<nblocks, nthreads>>>(inputimage, outputbuffer, blockmap);
/* after */}

__global__
kernel(byte* inputeimage, byte* outputbuffer, int* blockmap) {
    __shared__ int *skipFlag;


    if((blockDim.x * threadIdx.y + threadIdx.x) == 0)
    {
        *skipFlag = g_bMap[blockIdx.y * gridDim.x + blockIdx.x];
    }

    if(*skipFlag == 0)
    {
              /* recursive job */
    }
}
4

2 に答える 2

1

はい、これを行うことができますが、表示されたカーネル コードは正確にそれを行う方法ではありません。ブロックごとに整数フラグが必要だと仮定すると、コードは次のようになります。

__global__
kernel(byte* inputeimage, byte* outputbuffer, int* blockmap) {
    __shared__ int skipFlag;


    if (threadIdx.x == 0)
    {
        skipFlag = g_bMap[blockIdx.x];
    }
    __syncthreads();

    if(skipFlag == 0)
    {
              /* recursive job */
    }
}

ここで、各ブロックの最初のスレッドは、その特定のブロックのフラグをグローバルからロードし、それを共有メモリ整数変数に格納します。ブロックごとの同期の後、各スレッドはその値を読み取り、それに応じて動作できます。

于 2013-03-15T16:16:08.023 に答える
0

私はあなたの質問をよく理解していませんでしたが、その方法は正しいようです。skipFlag各ブロックは共有メモリに一意のものを持ち、取得したブロックはtrue残りのコードを実行しません。

そして、おそらく__syncthreads()両方のifの間が良い考えでしょう。

于 2013-03-15T16:06:07.960 に答える