CUDA :Когда использовать разделяемую память, а когда полагаться на кэширование L1?

После выпуска Compute Capability 2.0 (Fermi )я задался вопросом, остались ли еще варианты использования разделяемой памяти. То есть, когда лучше использовать разделяемую память, чем просто позволить L1 выполнять свою магию в фоновом режиме?

Разделяемая память предназначена только для того, чтобы алгоритмы, разработанные для CC < 2.0, могли эффективно работать без модификаций?

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

При использовании разделяемой памяти гарантируется, что значение, которое было помещено туда, останется там на протяжении всего блока. Это в отличие от значений в L1, которые вытесняются, если они не используются достаточно часто. Есть ли случаи, когда такие редко используемые данные лучше кэшировать в общей памяти, чем позволить L1 управлять ими на основе шаблона использования, который на самом деле имеет алгоритм?