Почему постоянный объем памяти ограничен в CUDA?
Согласно "Руководство по программированию CUDA C" , постоянный доступ к памяти выигрывает, только если используется кэш многопроцессорной константы (раздел 5.3.2.4) 1. В противном случае может быть даже больше запросов на память для полувращения, чем в случае чтения объединенной глобальной памяти. Итак, почему размер постоянной памяти ограничен 64 КБ?
Еще один вопрос, чтобы не спрашивать дважды. Насколько я понимаю, в архитектуре Fermi кеш текстуры объединен с кэшем L2. Использует ли использование текстуры смысл или чтение глобальной памяти кэшируется таким же образом?
1 Постоянная память (раздел 5.3.2.4)
Постоянное пространство памяти находится в памяти устройства и кэшируется в кеше констант, указанном в разделах F.3.1 и F.4.1.
Для устройств с вычислительной способностью 1.x запрос постоянной памяти для основы сначала разбивается на два запроса, по одному для каждого полувращения, которые выдаются независимо.
Затем запрос разбивается на столько отдельных запросов, что в исходном запросе есть разные адреса памяти, уменьшая пропускную способность на коэффициент, равный количеству отдельных запросов.
Полученные запросы затем обслуживаются при пропускной способности постоянного кеша в случае попадания в кэш или при пропускной способности памяти устройства в противном случае.
Ответы
Ответ 1
Размер постоянной памяти составляет 64 КБ для возможностей вычисления 1.0-3.0. Рабочий набор кэш-памяти составляет всего 8 КБ (см. Руководство по программированию CUDA v4.2 Таблица F-2).
Постоянная память используется драйвером, компилятором и переменными, объявленными __device__ __constant__. Драйвер использует постоянную память для связи параметров, привязок текстур и т.д. Компилятор использует константы во многих инструкциях (см. Разборку).
Переменные, помещенные в постоянную память, можно читать и записывать с использованием функций времени выполнения cudaMemcpyToSymbol() и cudaMemcpyFromSymbol() (см. раздел B.2.2 Руководства по программированию CUDA v4.2). Постоянная память находится в памяти устройства, но доступ к ней осуществляется через постоянный кеш.
В текстуре Ферми константа, L1 и I-Cache - это все кэши уровня 1 внутри или вокруг каждого SM. Все уровни 1 кэшируют память устройства доступа через кэш L2.
Постоянное ограничение 64 Кбайт - за CUmodule, который является модулем компиляции CUDA. Концепция CUmodule скрыта в среде выполнения CUDA, но доступна API-интерфейсом CUDA Driver.
