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

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

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

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

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

Ответ 1

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

Данные о кеше L1 не высылаются, если они не используются достаточно часто. Обычно он выдается, когда делается запрос для области памяти, которая ранее не была в кеше, и адрес которой разрешен только для того, который уже используется. Я не знаю точного алгоритма кэширования, используемого NVidia, но, предполагая регулярную ассоциативную ассоциацию n-way, тогда каждая запись в памяти может быть кэширована только в небольшом подмножестве всего кеша, на основе этого адреса

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

Ответ 2

Загрузка/хранение глобальной памяти подчиняется правилам коалесценции, даже если данные находятся в кеше, но разделяемая память намного более гибкая с точки зрения произвольного доступа. Я попытался использовать кеширование L1 для хранения/вычисления гистограммы, и это заканчивается тем, что намного, намного медленнее, чем использование разделяемой памяти из-за шаблона полуслучайного доступа.

Кроме того, согласно сотруднику NVIDIA, текущие кеши L1 являются сквозными (сразу же записываются в кеш L2), что замедлит вашу программу.

Итак, в основном, кэши мешают вам, если вам действительно нужна производительность.