Пиковая пропускная способность ядра cuda на графическом процессоре NVIDA - PullRequest
5 голосов
/ 06 августа 2011

У меня есть вопрос о пропускной способности ядра, работающего на графическом процессоре. Предполагая, что его занятость равна 0,5, размер блока равен 256: руководство по программированию утверждает, что лучше иметь много блоков, чтобы они могли скрыть задержку памяти и т. Д. Но я не понимаю, почему это правильно. Потому что как только ядро ​​получит количество деформаций на потоковый мультипроцессор = 24, то есть 3 блока, оно достигнет пиковой пропускной способности. Таким образом, наличие более 24 деформаций (или 3 блоков) ничего не изменит в пропускную способность.

Я что-то упустил? Кто-нибудь может меня поправить?

Ответы [ 2 ]

6 голосов
/ 08 августа 2011

Хотя верно, что SM с низкой загрузкой не могут в достаточной степени скрыть задержку, важно понимать это:

Более высокая занятость! = Более высокая пропускная способность!

Занятость - это просто мера того, сколько работы доступно SM для выбора в любой момент.Наличие большего количества резидентных деформаций дает SM больше возможностей для выполнения полезной работы, в то время как другие деформации ожидают результатов (результатов доступа к памяти или вычислений - оба имеют ненулевую задержку).

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

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

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

Вы не обязательно хотите иметь действительно крошечные блоки (поскольку SM на Fermi поддерживают не более 8 резидентных блоков), но иметь блоки из 128-512 потоков.часто более эффективен, чем использование блоков с 1024 потоками.

1 голос
/ 07 августа 2011

Наличие более 3 блоков не изменит пропускную способность в вашем случае, если у вас есть только один SM на вашей карте с поддержкой cuda.Обычно у вас есть 8 или более SM в одном GPU.

Кроме того, количество блоков, которые будут работать на одном SM, зависит не только от количества перекосов.Это только один ограничивающий фактор, есть много других факторов. CUDA Occupancy Calculator - отличный инструмент, чтобы увидеть занятость вашего ядра.

...