Параллельная запись в том же месте глобальной памяти
У меня есть несколько блоков, каждый из которых имеет целые числа в массиве разделяемой памяти размером 512. Как я могу проверить, содержит ли массив в каждом блоке нуль в качестве элемента?
То, что я делаю, это создание массива, который находится в глобальной памяти. Размер этого массива зависит от количества блоков, и он инициализируется равным 0. Следовательно, каждый блок записывается в a[blockid] = 1, если массив разделяемой памяти содержит нуль.
Моя проблема в том, что у меня одновременно есть несколько потоков в одном блоке. То есть, если массив в общей памяти содержит более одного нуля, то несколько потоков будут писать a[blockid] = 1. Это породит любую проблему?
Другими словами, было бы проблемой, если 2 потока записывают одно и то же значение в тот же самый элемент массива в глобальной памяти?
Ответы
Ответ 1
В модели исполнения CUDA нет гарантий, что каждая одновременная запись из потоков одного и того же блока в одно и то же место глобальной памяти будет успешной. По крайней мере одна запись будет работать, но модель программирования не гарантирует, сколько будет выполняться транзакций записи, или в каком порядке они появятся, если будет выполнено несколько транзакций.
Если это проблема, то лучший подход (с точки зрения правильности) должен состоять в том, чтобы только один поток из каждого блока выполнял глобальную запись. Вы можете либо использовать флаг общей памяти, установленный атомарно, либо операцию сокращения, чтобы определить, должно ли оно быть задано. Который вы выберете, может зависеть от того, сколько нулей, вероятно, будет. Чем больше нулей есть, тем более привлекательным будет сокращение. CUDA включает в себя операторы уровня __any() и __all(), которые могут быть встроены в очень эффективное булево сокращение в нескольких строках кода.
Ответ 2
Для программы CUDA, если несколько потоков в warp записываются в одно и то же место, то местоположение будет обновляться, но оно undefined, сколько раз обновление местоположения (то есть сколько фактических записи происходят последовательно), и это undefined, который поток будет писать последним (т.е. какой поток выиграет гонку).
Для устройств вычислительной возможности 2.x, если несколько потоков в warp записываются на один и тот же адрес, тогда только один поток фактически выполнит запись, этот поток - undefined.
Из Руководство по программированию CUDA C раздел F.4.2:
Если неатомическая команда, выполняемая warp, записывает одно и то же место в глобальной памяти для более чем одного из нитей warp, только один поток выполняет запись, а какой поток - undefined.
Дополнительную информацию см. в разделе 4.1 руководства.
Другими словами, если все потоки, записывающие в заданное место, записывают одно и то же значение, тогда это безопасно.
Ответ 3
Да, это будет проблема, называемая Race Condition.
Вы должны рассмотреть synchronizing доступ к глобальным данным через process Semaphores
Ответ 4
В то время как не Mutex или семафор, CUDA содержит примитив синхронизации, который вы можете использовать для сериализации доступа к данному сегменту кода или ячейке памяти. С помощью функции __syncthreads() вы можете создать барьер, чтобы любые поточные блоки в точке вызова команды до тех пор, пока все потоки в данном блоке не выполнили команду __syncthreads(). Таким образом, вы можете надеяться на сериализацию доступа к вашему месту памяти и избежать ситуации, когда два потока должны записываться в одно и то же место памяти одновременно. Единственное предупреждение состоит в том, что все потоки должны в какой-то момент выполнить __syncthreads(), иначе вы столкнетесь с ситуацией с мертвой блокировкой. Поэтому не помещайте вызов внутри некоторого условного if-оператора, где некоторые потоки никогда не смогут выполнить команду. Если вы подходите к своей проблеме вроде этого, необходимо будет предусмотреть некоторые потоки, которые изначально не вызывали __syncthreads(), чтобы вызвать функцию позже, чтобы избежать тупиковой ситуации.
