Как иметь атомную нагрузку в CUDA

Question

Мой вопрос, как я могу иметь атомную нагрузку в CUDA. Атомный обмен может подражать атомному хранилищу. Может ли атомная нагрузка недорогим образом эмулироваться аналогичным образом? Я могу использовать атомарное добавление с 0 для атомарной загрузки контента, но я думаю, что это дорого, потому что он выполняет атомарное чтение-изменение-запись вместо только чтения.

Answer 1

В дополнение к использованию volatile как рекомендовано в другом ответе, используя __threadfence соответственно также требуется получить атомную загрузку с безопасным упорядочением памяти.

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

Обычное чтение может повторно использовать предыдущую загрузку, которая уже есть в регистре, и, таким образом, может не отражать изменения, сделанные другими SM с желаемым упорядочением памяти. Например, int *flag = ...; while (*flag) { ... } может только читать flag один раз и повторно использовать это значение для каждой итерации цикла. Если вы ожидаете, что другой поток изменит значение флага, вы никогда не увидите изменения. volatile Модификатор гарантирует, что значение фактически читается из памяти при каждом доступе. См. Документацию CUDA для volatile для получения дополнительной информации.

Кроме того, вам нужно использовать ограничитель памяти для обеспечения правильного упорядочения памяти в вызывающем потоке. Без ограничений вы получаете "расслабленную" семантику на языке C++11, и это может быть небезопасно при использовании атомарного взаимодействия.

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

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

Правильно расположенные ограждения памяти предотвращают подобные проблемы, обеспечивая, чтобы переупорядочение команд не влияло на желаемое упорядочение памяти и чтобы предыдущие записи были видны другим потокам. __threadfence() и друзья также включены в документы CUDA.

Собрав все это вместе, написание собственного метода атомарной загрузки в CUDA выглядит примерно так:

// addr must be aligned properly.
__device__ unsigned int atomicLoad(const unsigned int *addr)
{
  const volatile unsigned int *vaddr = addr; // volatile to bypass cache
  const unsigned int value = *vaddr;
  // fence to ensure that dependent reads are correctly ordered
  __threadfence(); 
  return value; 
}

// addr must be aligned properly.
__device__ void atomicStore(unsigned int *addr, unsigned int value)
{
  volatile unsigned int *vaddr = addr; // volatile to bypass cache
  // fence to ensure that previous non-atomic stores are visible to other threads
  __threadfence(); 
  *vaddr = value;
}

Это можно записать аналогично для других размеров без нагрузки / хранения.

Из разговоров с некоторыми разработчиками NVIDIA, которые работают над атомарностью CUDA, похоже, что мы должны начать видеть лучшую поддержку атомарности в CUDA, и PTX уже содержит инструкции загрузки / хранения с семантикой упорядочения памяти получения / выпуска - но нет никакого способа доступ к ним в настоящее время, не прибегая к встроенному PTX. Они надеются добавить их в этом году. Как только они на месте, полный std::atomic реализация не должна сильно отставать.

Answer 2

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

Есть две квази- альтернативы, с их преимуществами и недостатками:

1. Используйте неоперативное атомарное чтение-изменение-запись, как вы предлагаете. Я дал аналогичный ответ в прошлом. Гарантированная атомарность и согласованность памяти, но вы платите за ненужную запись.

2. На практике, вторая самая близкая вещь к атомной нагрузке может быть маркировка переменной volatileхотя, строго говоря, семантика совершенно иная. Язык не гарантирует атомарность нагрузки (например, теоретически вы можете получить разорванное чтение), но вы гарантированно получите самое актуальное значение. Но на практике, как указано в комментариях @Robert Crovella, невозможно получить разорванное чтение для правильно выровненных транзакций длиной не более 32 байт, что делает их атомарными.

Решение 2 довольно хакерское, и я не рекомендую его, но в настоящее время оно является единственной альтернативой 1. Без записи. Идеальным решением было бы добавить способ выражать атомарные нагрузки непосредственно в языке.