Общая память cuda и планирование выполнения блоков

Question

Я хотел бы прояснить состояние выполнения с разделяемой памятью CUDA и выполнить блок на основе количества общей памяти, используемой на блок.

государственный

Я использую GTX480 NVIDIA карту, которая имеет 48 КБ общей памяти на блок и 15 потоковых мультипроцессоров. Итак, если я объявляю ядро ​​с 15 блоками, каждое использует 48 КБ разделяемой памяти, и никаких других ограничений не достигается (регистры, максимальное количество потоков на блок и т. Д.), Каждый блок работает до одного SM(из 15) до конца. В этом случае требуется только планирование между перекосами одного и того же блока.

Вопрос

Итак, мой сценарий недопонимания:
Я называю ядро ​​с 30 блоками, так что 2 блока находятся на каждом SM. Теперь планировщику на каждом СМ приходится иметь дело с перекосами из разных блоков. Но только когда один блок заканчивает свое выполнение, деформации другого блока выполняются на SM из-за использования общего объема общей памяти (48 КБ на SM). Если этого не произойдет, и если различные блоки планируют выполнение на одном и том же SM, результат может быть неправильным, поскольку один блок может считывать значения, загруженные из другого, в общую память. Я прав?

Answer 1

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

---

Кстати, из соображений производительности обычно лучше иметь как минимум два блока, работающих на SM, потому что

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

Конечно, могут быть причины, по которым увеличение общей памяти на блок дает большее ускорение, чем при работе с несколькими блоками на SM.