atomicInc() не работает

Question

Я попробовал ниже программу, используя atomicInc().

__global__ void ker(int *count)
{
    int n=1;
    int x = atomicInc ((unsigned int *)&count[0],n);
    CUPRINTF("In kernel count is %d\n",count[0]);
}

int main()
{
    int hitCount[1];
    int *hitCount_d;

    hitCount[0]=1;
    cudaMalloc((void **)&hitCount_d,1*sizeof(int));

    cudaMemcpy(&hitCount_d[0],&hitCount[0],1*sizeof(int),cudaMemcpyHostToDevice);

    ker<<<1,4>>>(hitCount_d);

    cudaMemcpy(&hitCount[0],&hitCount_d[0],1*sizeof(int),cudaMemcpyDeviceToHost);

    printf("count is %d\n",hitCount[0]);
  return 0;
}

Выход:

In kernel count is 1
In kernel count is 1
In kernel count is 1
In kernel count is 1

count is 1

Я не понимаю, почему это не увеличивается. Может ли кто-нибудь помочь

Answer 1

Ссылаясь на документацию, atomicInc Является ли это:

для следующих:

atomicInc ((unsigned int *)&count[0],n);

вычислить:

((count[0] >= n) ? 0 : (count[0]+1))

и сохранить результат обратно в count[0]

(Если вы не уверены, что ? Оператор, посмотри здесь)

Так как вы прошли n = 1 и count[0] начинается в 1, atomicInc никогда не увеличивает переменную count[0] за 1.

Если вы хотите, чтобы его значение превышало 1, передайте большее значение для n,

Переменная n фактически действует как "значение ролловера" для процесса приращения. Когда переменная, которая будет увеличена, фактически достигает значения n, следующий atomicInc сбросит его на ноль.

Хотя вы не задавали вопрос, вы можете спросить: "Почему я никогда не вижу нулевое значение, если я нажимаю на значение прокрутки?"

Чтобы ответить на это, вы должны помнить, что все 4 ваших потока выполняются в режиме lockstep. Все 4 из них выполняют atomicInc Инструкция перед выполнением любого последующего оператора печати.

Поэтому у нас есть переменная count[0] который начинается в 1.

1. Первый поток для выполнения атомарного сбрасывает его в ноль.
2. Следующий поток увеличивает его до 1.
3. Третий поток сбрасывает его на ноль.
4. Четвертый и последний поток увеличивает его до 1.

Затем все 4 потока выводят значение.

В качестве другого эксперимента попробуйте запустить 5 потоков вместо 4, посмотрите, сможете ли вы предсказать, какое значение будет напечатано.

ker<<<1,5>>>(hitCount_d);

Как указано @talonmies в комментариях, если вы поменяете местами atomicInc для atomicAdd:

int x = atomicAdd ((unsigned int *)&count[0],n);

Вы получите результаты, которые вы, вероятно, ожидали.