Cuda atomics меняет флаг

Question

У меня есть кусок последовательного кода, который делает что-то вроде этого

if( ! variable )
{
  do some initialization here 
  variable = true;
}

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

Answer 1

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

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

Атомика сама по себе может использоваться только для очень ограниченной, в основном единственной операции, над одной переменной. Но атомика может быть использована для создания критической секции.

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

__syncthreads();
if (threadIdx.x == 0)
  acquire_semaphore(&sem);
__syncthreads();
  //begin critical section
  // ... your critical section code goes here
  //end critical section
__syncthreads();
if (threadIdx.x == 0)
  release_semaphore(&sem);
__syncthreads();

Перед ядром определите эти вспомогательные функции и переменную устройства:

__device__ volatile int sem = 0;

__device__ void acquire_semaphore(volatile int *lock){
  while (atomicCAS((int *)lock, 0, 1) != 0);
  }

__device__ void release_semaphore(volatile int *lock){
  *lock = 0;
  __threadfence();
  }

Я протестировал и успешно использовал приведенный выше код. Обратите внимание, что он по существу осуществляет арбитраж между блоками потоков, используя поток 0 в каждом блоке потоков в качестве запросчика. Вы должны далее условие (например, if (threadIdx.x < ...)) код критической секции, если вы хотите, чтобы только один поток в блоке-победителе выполнял код критической секции.

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

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

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