Действительно ли атомарные функции делают переменные переменными в CUDA?

Question

Я написал очень простой код, запрашивающий поток 0 обновить глобальную переменную, в то время как другие потоки продолжают читать эту переменную. Но я обнаружил, что другие потоки действительно не получают значение.

Код здесь, это довольно просто. Кто-нибудь может дать мне какое-нибудь предложение, как это исправить? большое спасибо

__global__ void addKernel(int *c)
{
int i = threadIdx.x;
int j = 0;
if (i == 0)
{
    while(*c < 2000){
        int temp = *c;
        printf("*c = %d\n",*c);
        atomicCAS(c,temp, temp+1);
    }       
}else{
    while(*c < 1000)
    {
        j++;
    }
}

}

Answer 1

Я хотел бы привести аналогию: представьте на секунду, что атомарные операции являются мьютексами: чтобы программа была четко определена, два потока, обращающиеся к общему ресурсу, должны оба согласиться использовать мьютекс для доступа к ресурсу исключительно. Если один из потоков обращается к ресурсу без предварительного удержания мьютекса, результат не определен.

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

Другими словами, это:

atomicCAS(c,temp, temp+1);

Содержит атомную загрузку-сравнение-магазин. Полученная инструкция будет загружена в глобальную память c Выполните сравнение и пройдите до глобальной памяти, чтобы сохранить новое значение.

Но это:

while(*c < 2000)

Не является атомным в любом случае. Компилятор (и аппаратное обеспечение) понятия не имеет, что c возможно, был изменен другим потоком. Таким образом, вместо того, чтобы перейти к глобальной памяти, он будет просто читать из самого быстрого доступного кеша. Возможно, компилятор даже поместит переменную в регистр, потому что он не увидит, чтобы кто-то еще модифицировал ее в текущем потоке.

То, что вы хотели бы, это что-то вроде (воображаемое):

while (atomicLoad(c) < 2000)

Но, насколько мне известно, на момент написания этой статьи в CUDA такой конструкции не было.

В связи с этим volatile Может помочь квалификатор: он говорит компилятору не оптимизировать переменную и считает ее "модифицируемой из внешних источников". Это вызовет загрузку для каждого чтения переменной, хотя я не уверен, что эта загрузка обходит все кэши. На практике это может сработать, но в теории я не думаю, что вы должны полагаться на это. Кроме того, это также отключит любую оптимизацию для этой переменной (например, постоянное распространение или продвижение переменной в регистр для лучшей производительности).

Вы можете попробовать следующий хак (я не пробовал):

while(atomicAdd(c, 0) < 2000)

Это выдаст атомарную инструкцию, которая загружается из глобальной памяти, и поэтому должна увидеть самое последнее значение c, Тем не менее, он также вводит (бесполезный в этом случае) атомарный магазин.