c - 定数 for ループが間違った結果をもたらす

Question

CUDA デバイス用のヒストグラム カーネルを作成しようとしています。これはNVIDIA の論文に基づいています。

アイデアは、すべてのスレッドが特定の部分 (私の場合はボリューム) の部分ヒストグラムを計算し、それを共有メモリのブロックに書き込むというものです。しかし、アルゴリズムで奇妙な問題が発生したため、カーネルを重要な部分に分解しました。

__global__ void calcHist64()
{
    extern __shared__ unsigned char partialHistograms[];

    //a unique sequential thread id within this block, used to determine the memory in which to write the partial histogram
    unsigned int seqTid = threadIdx.x + threadIdx.y * blockDim.x;       
#pragma unroll
    for(int i = 0; i < 255; ++i)
    {
        //increment the thread's partial histogram value
        partialHistograms[seqTid]++; 
    }
    //each partial histogram should now be 255
    //Output the value for every thread in a certain block
    if(blockIdx.x == 0 && blockIdx.y == 31)
        printf("Partial[%i][%i]: %i\n", threadIdx.x, threadIdx.y, partialHistograms[partialHistRoot]);  
}

カーネルは次の方法で呼び出されます。

int sharedMemory = 4096;
dim blocks(32, 32, 1);
dim3 threadsPerBlock(8,8,1);
calcHist64<<<blocks, threadsPerBlock, sharedMemory>>>();

blockIdx.x各部分ヒストグラムの値は 255 であると予想しています。ただし、これは最初の数ブロック (低/ )にのみ当てはまりますblockIdx.y。他のブロックの値は大きく異なります。最後のブロック ( blockIdx.y == 31) の値は 239 または 240 です。

この振る舞いを説明することはできません。結局、ちょうど 255 回実行される for ループの定数です。各スレッドは共有メモリの異なる部分にアクセスするため、競合状態は発生しません。

誰でもこの動作を説明できますか?