Почему постоянный объем памяти ограничен в CUDA?

Question

Согласно "CUDA C Programming Guide", постоянный доступ к памяти может быть полезен только при попадании кеша многопроцессорной константы (раздел 5.3.2.4) . В противном случае может быть даже больше запросов на память для полувращения, чем в случае чтения объединенной глобальной памяти. Итак, почему размер постоянной памяти ограничен 64 КБ?

Еще один вопрос, чтобы не спрашивать дважды. Насколько я понимаю, в архитектуре Fermi кеш текстуры объединен с кэшем L2. Сохраняется ли использование текстуры по-прежнему, или чтение глобальной памяти кэшируется таким же образом?

---

Постоянная память (раздел 5.3.2.4)

Постоянное пространство памяти находится в памяти устройства и кэшируется в постоянном кэше, указанного в разделах F.3.1 и F. 4.1.

Для устройств с вычислительной способностью 1.x запрос постоянной памяти для основы сначала разбивается на два запроса, по одному для каждого полувращения, которые выдаются независимо.

Затем запрос разбивается на столько раздельных запросов, что в исходном запросе есть разные адреса памяти, что уменьшает пропускную способность на коэффициент, равный количеству отдельных запросов.

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

Answer 1

Размер постоянной памяти составляет 64 КБ для вычислительной способности 1.0-3.0. Рабочий набор кэш-памяти составляет всего 8 КБ (см. Руководство по программированию CUDA v4.2 Таблица F-2).

Постоянная память используется указанным драйвером, компилятором и переменными __device__ __constant__. Драйвер использует постоянную память для связи параметров, привязки текстур и т. Д. Компилятор использует константы во многих инструкциях (см. Разборку).

Переменные, размещенные в постоянной памяти, могут быть считаны и записаны с использованием функций времени выполнения хоста cudaMemcpyToSymbol() и cudaMemcpyFromSymbol() (см. Раздел B.2.2 Руководства по программированию CUDA v4.2). Постоянная память находится в памяти устройства, но доступ к ней осуществляется через постоянный кеш.

Текстура Fermi, постоянная, L1 и I-Cache - это все кеши уровня 1 внутри или вокруг каждого SM. Все уровни 1 кэшируют память устройства доступа через кэш L2.

Постоянный предел в 64 Кбайт за CUmodule, который является модулем компиляции CUDA. Концепция CUmodule скрыта в среде исполнения CUDA, но доступна API-интерфейсом CUDA.