Атомные операции в CUDA

Question

Следующая программа использовала реализацию атомарных блокировок из "Cuda By Example", но при запуске программы моя машина зависла. Может кто-нибудь сказать мне, что не так с моей программой? большое спасибо

Ифэй

#include <stdio.h>


__global__ void test()
{
    __shared__ int i, mutex;

    if (threadIdx.x == 0) {
       i = 0;
       mutex = 0;
    }
    __syncthreads();

    while( atomicCAS(&mutex, 0, 1) != 0);
    i++;
    printf("thread %d: %d\n", threadIdx.x, i);
    atomicExch(&mutex,0);
}

Answer 1

Вот теория. Я надеюсь, что вы знакомы с концепцией деформации. В цикле while все потоки внутри деформации войдут в цикл while. Выходит только один, а остальные потоки будут находиться внутри цикла while. Это введет расходящуюся ветвь, делающую поток, который вышел из цикла while, останавливается, пока ветвь не сходится снова. Поскольку этот поток является единственным, который может освободить мьютекс, этого никогда не произойдет, потому что он ожидает, пока другие потоки не сойдутся.