Играюсь тут с CUDA, хочу сделать параллельную гистограмму, в перспективе вообще без атомиков.

__global__
void histo(const unsigned int* const vals,    //INPUT
                 unsigned int* const histo,   //OUPUT
                          int        numVals)
{
    extern __shared__ unsigned int sdata[];

    const unsigned int tid        = threadIdx.x;
    const unsigned int blockId    = blockDim.x * blockIdx.x;
    const unsigned int myId       = tid + blockId;

    if (myId > numVals) {
        return;
    }

    sdata[tid] = 0;
    __syncthreads();

    const unsigned int curVal     = vals[myId];
    if (tid == curVal) {
        atomicAdd(&(sdata[curVal]), 1);
    }

    __syncthreads();
    const unsigned int blockHistoVal = sdata[tid];
    if (blockHistoVal != 0) {
        atomicAdd(&(histo[tid]), blockHistoVal);
    }
}

histo<<<numBlocksForElements, MAX_THREADS, MAX_THREADS*sizeof(unsigned int)>>> (d_vals, d_histo, numElems);

То есть идея такая. Запускаю я этот kernel с 1024 потоками, с соответствующими параметрами запуска. Предполагаю, что это будет сгруппировано по блокам с 1024 потоками. Соответственно для каждого блока я соответствующие очищаю shared memory, синхронизую, делаю пока просто атомиком тривиальную гистограмму над shared memory, синхронизую, и затем атомиком уже из всех shared memory агрегирую в результат. Знаю, что это все тормознуто, но дело не в этом. Какого то черта здест гонки явные возникают, а вот почему — понять не могу. Предполагаю, что __syncthreads() синхронизует потоки внутри thread block с которым связана shared memory, соответственно ну никак не ожидаю я здесь гонок. А они есть. Явно я что то не понимаю в устройстве как это все работает.