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

Question

Я не могу понять это сам, что является лучшим способом обеспечить постоянную память, используемую в моем ядре. Есть похожий вопрос на http://stackru...r-pleasant-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, Два не эквивалентны, второй гарантирует, что он остается неизменным после инициализации. Если вам нужна динамика, то вам нужно использовать что-то похожее на первый способ.