Постоянное использование памяти в коде CUDA

Question

Сам не могу понять, как лучше всего обеспечить постоянную память, используемую в моем ядре.Есть похожий вопрос на http://stackoverflow. .. r-приятно-way .Я работаю с GTX580 и собираю только для возможности 2.0.Мое ядро ​​выглядит так:

__global__ Foo(const int *src, float *result) {...}

Я выполняю следующий код на хосте:

cudaMalloc(src, size);
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice);
Foo<<<...>>>(src, result);

. Альтернативный способ - добавить

__constant__ src[size];

в файл .cu,удалить указатель src из ядра и выполнить

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice);
Foo<<<...>>>(result);

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

Answer 1

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

Типичными вещами, которые могут подходить для памяти __constant__, могут быть коэффициенты модели, веса и другие постоянные значения, которые необходимо установить во время выполнения.Например, на графических процессорах Fermi список аргументов ядра хранится в постоянной памяти.Но если доступ к содержимому осуществляется неравномерно, а тип или размер членов не постоянен от вызова к вызову, тогда предпочтительна обычная глобальная память.

Также имейте в виду, что существует ограничение в 64 КБ постоянной памяти на контекст GPU, поэтому нецелесообразно хранить очень большие объемы данных в постоянной памяти.Если вам нужно много доступного только для чтения хранилища с кешем, возможно, стоит попробовать привязать данные к текстуре и посмотреть, какова производительность.На картах pre-Fermi это обычно дает удобный выигрыш в производительности, на Fermi результаты могут быть менее предсказуемыми по сравнению с глобальной памятью из-за улучшенной структуры кэша в этой архитектуре.

Answer 2

Первый метод гарантирует постоянную память внутри функции Foo. Два не эквивалентны, второй гарантирует, что он остается неизменным после инициализации. Если вам нужна динамика, то вам нужно использовать что-то похожее на первый способ.