Как эффективно собирать данные из потоков в CUDA?

Question

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

Я прохожу мимо огромного массива с достаточным пространством для всех потоков, хранящим 4 решения (4 двухместных для каждого решения), а другой - с количеством решений на поток, однако это наивное решение и является текущим узким местом моего ядра.

Я действительно хочу оптимизировать эту тему. Основная проблема заключается в конкатенации переменной количества решений на поток в одном массиве.

Answer 1

Функциональность, которую вы ищете, называется уплотнением потока.

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

Однако, учитывая более распространенную ситуацию, когда имеется множество результатов, лучшим решением будет выполнение операции уплотнения в виде отдельного шага. Один из способов сделать это - thrust::copy_if. См. this question для получения дополнительной информации.