CUDA AtomicAdd быстрее, чем небезопасный доступ

Question

Я делаю свои первые шаги с CUDA и пытаюсь реализовать простой расчет гистограммы.Я просто делю массив на несколько блоков и потоков и, при необходимости, позволяю потоку вычислять две итерации цикла while.

Это мое ядро ​​CUDA, и первое, что я хотел увидеть, - насколько медленнее atomicAddдоступ сравнивается с небезопасной записью в массив гистограмм.Но когда я проверил вывод nvprof, вывод показывает меня для массива из 1073741824 элементов (1<<30) и BlockSize из 128 и GridSize из 8388608 этого вывода:

13.49%  126.33ms     1  126.33ms  126.33ms  126.33ms  histogram_unsafe(int*, int*)
12.01%  112.49ms     1  112.49ms  112.49ms  112.49ms  histogram_safe(int*, int*)

(При такой конфигурации каждый поток выполняет одну итерацию цикла)

Но в моей фантазии небезопасный код должен работать намного быстрее?По крайней мере, небезопасный код не дает правильного результата, в то время как безопасный код делает.Но я не понимаю, почему небезопасная версия медленнее?Есть ли какое-то эмпирическое правило о том, что такое атмо с CUDA, чего я не знаю?

Вот мой код ядра:

__global__ void histogram_[unsafe/save]( int *input, int *hist ) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    while (tid < N) {
        int bin = input[tid] & MASK;

        // Of course I use only one version at a time
        hist[bin]++; // unsafe
        atomicAdd(&hist[bin], 1); // safe

        tid += blockDim.x * gridDim.x;
    }
}