#c #c #cuda
#c #c #cuda
Вопрос:
У меня есть несколько блоков, каждый из которых содержит несколько целых чисел в массиве общей памяти размером 512. Как я могу проверить, содержит ли массив в каждом блоке ноль в качестве элемента?
Что я делаю, так это создаю массив, который находится в глобальной памяти. Размер этого массива зависит от количества блоков, и он инициализируется значением 0. Следовательно, каждый блок выполняет запись в a[blockid] = 1 , если массив общей памяти содержит ноль.
Моя проблема заключается в том, что у меня одновременно записывается несколько потоков в одном блоке. То есть, если массив в общей памяти содержит более одного нуля, то запись будет производиться несколькими потоками a[blockid] = 1 . Вызовет ли это какие-либо проблемы?
Другими словами, будет ли проблемой, если 2 потока запишут одно и то же значение в один и тот же элемент массива в глобальной памяти?
Ответ №1:
Для программы CUDA, если несколько потоков в warp выполняют запись в одно и то же местоположение, то это местоположение будет обновлено, но не определено сколько раз это местоположение обновляется (т. Е. Сколько фактических записей происходит последовательно), и не определено какой поток будет записывать последним (т. Е. Какой поток выиграет гонку).
Для устройств с вычислительными возможностями 2.x, если несколько потоков в warp записывают на один и тот же адрес, тогда только один поток фактически выполнит запись, какой поток не определен.
Из раздела F.4.2 руководства по программированию на CUDA C:
Если неатомная инструкция, выполняемая warp, выполняет запись в одну и ту же ячейку глобальной памяти для более чем одного из потоков warp, запись выполняет только один поток, и какой поток это делает, не определено.
Смотрите также раздел 4.1 руководства для получения дополнительной информации.
Другими словами, если все потоки, выполняющие запись в заданную ячейку, записывают одно и то же значение, то это безопасно.
Ответ №2:
В модели выполнения CUDA нет никаких гарантий, что каждая одновременная запись из потоков в одном блоке в одну и ту же глобальную ячейку памяти будет успешной. По крайней мере, одна запись будет работать, но модель программирования не гарантирует, сколько транзакций записи произойдет или в каком порядке они будут выполняться, если выполняется более одной транзакции.
Если это проблема, то лучшим подходом (с точки зрения корректности) было бы, чтобы глобальную запись выполнял только один поток из каждого блока. Вы можете либо использовать флаг общей памяти, установленный атомарно, либо операцию сокращения, чтобы определить, следует ли устанавливать значение. То, что вы выберете, может зависеть от того, сколько нулей там может быть. Чем больше нулей, тем привлекательнее будет сокращение. CUDA включает в себя операторы уровня деформации __any() и __all() , которые могут быть встроены в очень эффективное логическое сокращение в нескольких строках кода.
Комментарии:
1. Мой 1 за ответ с точки зрения CUDA, который ищет OP, а не с точки зрения среды разработки C / C .
2. Пожалуйста, смотрите мой ответ для получения более подробной информации (не могу публиковать ссылки и цитаты в комментарии!). CUDA гарантирует, что если несколько потоков в warp выполняют запись в одно и то же местоположение, то по крайней мере одному потоку удастся выполнить запись в это местоположение, но какой это поток (или какой поток идет последним) не определен.
Ответ №3:
Да, это будет проблема, называемая как Race Condition .
Вам следует рассмотреть synchronizing возможность доступа к глобальным данным через process Semaphores
Комментарии:
1. ммм i c было бы лучше, если бы я использовал атомарную операцию или использовал алгоритм сокращения, чтобы проверить, содержит ли массив нуль?
2. @lina: Это зависит от того, можете ли вы сделать эти операции атомарными.. Синхронизация может быть простой, если вы немного ознакомитесь с основами… academictutorials.com/ipc/ipc-process-synchronization.asp hth
3. @Als: В CUDA отсутствуют глобальные механизмы синхронизации, такие как семафоры и мьютексы, поэтому, хотя в предложении нет ничего плохого, в данном случае оно неприменимо.
4. @talonmies: Я думаю, вы правы, мой ответ был с точки зрения C / C , и я должен признать, что я не разбираюсь в Cuda. Так получилось, что я ответил на это с предварительной точки зрения C / C .
5. @lina: Мои извинения, мой ответ, похоже, не подходит для Cuda, у меня нет опыта в этом. Я ответил на это с точки зрения C / C , и это было бы справедливо для C / C . В любом случае, я думаю, что ответ @ talomies кажется подходящим для вашего вопроса, учитывая среду разработки Cuda.
Ответ №4:
CUDA, не являясь мьютексом или семафором, содержит элемент синхронизации, который вы можете использовать для сериализации доступа к заданному сегменту кода или ячейке памяти. С помощью __syncthreads() функции вы можете создать барьер, чтобы любой данный поток блокировался в момент вызова команды до тех пор, пока все потоки в данном блоке не выполнят __syncthreads() команду. Надеюсь, таким образом вы сможете сериализовать доступ к вашей ячейке памяти и избежать ситуации, когда двум потокам необходимо выполнять запись в одну и ту же ячейку памяти одновременно. Единственное предупреждение заключается в том, что все потоки должны в какой-то момент выполняться __syncthreads() , иначе вы окажетесь в ситуации тупиковой блокировки. Поэтому не размещайте вызов внутри некоторого условного оператора if, где некоторые потоки могут никогда не выполнить команду. Если вы подходите к своей проблеме подобным образом, необходимо будет предусмотреть, чтобы потоки, которые изначально не вызываются __syncthreads() , вызывали функцию позже, чтобы избежать взаимоблокировки.