Одновременная запись в одну и ту же область глобальной памяти

Question

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

Я занимаюсь созданием массива, который находится в глобальной памяти.Размер этого массива зависит от количества блоков, и он инициализируется равным 0. Следовательно, каждый блок записывает в a[blockid] = 1, если массив разделяемой памяти содержит ноль.

Моя проблема в том, что у меня одновременно несколько потоков в одном блоке.То есть, если массив в разделяемой памяти содержит более одного нуля, то несколько потоков будут записывать a[blockid] = 1.Будет ли это вызвать какие-либо проблемы?

Другими словами, будет ли проблемой, если 2 потока записывают одинаковое значение в один и тот же элемент массива в глобальной памяти?

Answer 1

Для программы CUDA, если несколько потоков в деформации записывают в одно и то же местоположение, тогда местоположение будет обновлено, но оно не определено сколько раз местоположение обновляется (т. е. сколько фактических записей происходит последовательно), и undefined какой поток будет писать последним (т. е. какой поток выиграет гонку).

Для устройств с вычислительной возможностью 2.x, если несколько потоков в деформации записывают по одному и тому же адресу, тогда только один поток фактически выполняет запись, , который поток не определен.

Из руководства по программированию CUDA C , раздел F.4.2:

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

См. Также раздел 4.1 руководства для получения дополнительной информации.

Другими словами, если все потоки, записывающие в заданное местоположение, записывают одно и то же значение, тогда это безопасно.

Answer 2

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

Если это проблема, то лучшим подходом (с точки зрения корректности) было бы, чтобы только один поток из каждого блока выполнял глобальную запись.Вы можете использовать флаг общей памяти, установленный атомарно, или операцию сокращения, чтобы определить, следует ли установить значение.То, что вы выберете, может зависеть от того, сколько нулей может быть.Чем больше нулей, тем привлекательнее будет сокращение.CUDA включает в себя операторы уровня деформации __any() и __all(), которые могут быть встроены в очень эффективное логическое сокращение нескольких строк кода.

Answer 3

Да, это будет проблема, которая называется Race Condition.
Вам следует рассмотреть возможность synchronizing доступа к глобальным данным через process Semaphores

Answer 4

Хотя CUDA не является мьютексом или семафором, он содержит примитив синхронизации, который можно использовать для сериализации доступа к заданному сегменту кода или ячейке памяти.С помощью функции __syncthreads() вы можете создать барьер так, чтобы любой заданный поток блокировался в точке вызова команды, пока все потоки в данном блоке не выполнят команду __syncthreads().Таким образом, можно надеяться, что вы можете сериализовать доступ к вашей ячейке памяти и избежать ситуации, когда двум потокам необходимо одновременно выполнять запись в одну и ту же ячейку памяти.Единственное предупреждение состоит в том, что все потоки должны в какой-то момент выполнить __syncthreads(), иначе вы окажетесь в ситуации тупиковой блокировки.Поэтому не размещайте вызов внутри некоторого условного оператора if, где некоторые потоки могут никогда не выполнить команду.Если вы действительно подходите к своей проблеме, как это, необходимо будет предусмотреть некоторые условия для потоков, которые изначально не вызывают __syncthreads(), чтобы позже вызывать функцию, чтобы избежать тупика.