доступ к глобальной памяти для отдельных потоков

Question

Я пишу упрощенный raytracer. Идея состоит в том, что для каждого пикселя существует поток, который проходит определенную структуру (геометрию), которая находится в глобальной памяти.

Я вызываю свое ядро ​​так:

trace<<<gridDim, blockDim>>>(width, height, frameBuffer, scene)

Где scene - это структура, которая ранее была выделена с cudaMalloc. Каждый поток должен начать обход этой структуры, начиная с одного и того же узла, и есть вероятность, что многие параллельные потоки будут пытаться читать одни и те же узлы много раз. Означает ли это, что когда такие чтения происходят, это наносит вред степени параллелизма?

Учитывая, что геометрия велика, я бы предположил, что ее копирование не вариант. Я имею в виду, что вся обработка по-прежнему происходит довольно быстро, но мне было интересно, нужно ли с этим что-то делать или просто бросить на ветер.

Answer 1

Эффективный доступ к глобальной памяти из каждого потока зависит как от архитектуры вашего устройства, так и от вашего кода. Массивы, размещенные в глобальной памяти, выравниваются драйвером CUDA по 256-байтовым сегментам памяти. Устройство может получить доступ к глобальной памяти через 32-, 64- или 128-байтовые транзакции, которые выровнены по их размеру. Устройство объединяет глобальные загрузки памяти и сохраняет выпущенные потоки деформации в минимально возможное количество транзакций для минимизации пропускной способности DRAM. Неверно выровненный доступ к данным для устройств с вычислительной способностью менее 2,0 влияет на эффективную пропускную способность доступа к данным. Это не является серьезной проблемой при работе с устройством, которое имеет вычислительные возможности> 2,0. При этом, в значительной степени независимо от поколения вашего устройства, при доступе к глобальной памяти большими шагами эффективная пропускная способность становится низкой ( Reference ). Я предположил бы, что для произвольного доступа такое же поведение вероятно.

Answer 2

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

В соответствии со статьей здесь доступ к памяти может быть объединен, если имеется локальность данных и только в пределах перекосов.

Это означает, что если потоки внутри деформации пытаются получить доступ к областям памяти рядом друг с другом, они могут быть объединены. В вашем случае каждый поток пытается получить доступ к «одному и тому же» узлу, пока не встретит конечную точку, где они разветвляются.

Это означает, что доступ к памяти будет объединяться в рамках деформации, пока потоки не разветвляются.

Answer 3

Если вы не изменяете структуру во время чтения, что, как я полагаю, вы делаете (если вы, вероятно, визуализируете каждый кадр?), Тогда да, это снижает производительность и может привести к неопределенному поведению. Это называется состоянием гонки. Вы можете использовать атомарные операции, чтобы преодолеть этот тип проблемы. Использование атомарных операций гарантирует, что условия гонки не произойдут. Вы можете попытаться, вставив «сцену» в общую память, если сможете. Вы также можете попробовать использовать потоки для увеличения параллелизма, что также приводит к некоторой синхронизации с ядрами, работающими в одном потоке.