Хотя верно, что SM с низкой загрузкой не могут в достаточной степени скрыть задержку, важно понимать это:
Более высокая занятость! = Более высокая пропускная способность!
Занятость - это просто мера того, сколько работы доступно SM для выбора в любой момент.Наличие большего количества резидентных деформаций дает SM больше возможностей для выполнения полезной работы, в то время как другие деформации ожидают результатов (результатов доступа к памяти или вычислений - оба имеют ненулевую задержку).
Пропускная способность является мерой того, каквыполняется много работы в секунду, и хотя она может быть ограничена задержкой (и, следовательно, занятостью), она также может быть ограничена пропускной способностью памяти, пропускной способностью команд (количеством исполнительных блоков) и другими факторами.
Причина, по которой руководство по программированию утверждает, что лучше иметь несколько потоковых блоков, чем один большой блок потоков, заключается в том, что иногда лучше иметь возможность выполнять работу не только из других деформаций, но и из других блоков.Вот пример:
Представьте, что ваш блок большого потока должен загружать данные из глобальной памяти (с высокой задержкой) и сохранять их в общей памяти (с низкой задержкой), а затем должен немедленно выполнить __syncthreads()
.В этом случае, когда деформация завершает загрузку своих данных и запись их в разделяемую память, она должна затем ждать, пока все остальные потоки в блоке не сделают то же самое.Для большого блока это может занять довольно много времени.Но если SM занимает несколько блоков меньшего размера, SM может переключаться и выполнять работу из других блоков, ожидая, пока __syncthreads
будет удовлетворен в первом блоке.Это может помочь уменьшить время простоя графического процессора и повысить эффективность.
Вы не обязательно хотите иметь действительно крошечные блоки (поскольку SM на Fermi поддерживают не более 8 резидентных блоков), но иметь блоки из 128-512 потоков.часто более эффективен, чем использование блоков с 1024 потоками.