На этом шаге мы рассмотрим константную память.
Константная память выделяется непосредственно в коде программы при помощи спецификатора __constant__. Все нити сетки могут читать из нее данные, и чтение из нее кэшируется. CPU имеет доступ к ней как на чтение, так и на запись. Ниже приведены функции, при помощи которых CPU может взаимодействовать с константной памятью.
cudaError_t cudaMemcpyToSymbol(const char *symbol, const void *src, size_t count, size_t offset, enum cudaMemcpyKind kind); cudaError_t cudaMemcpyFromSymbol(void *dst, const char *symbol, size_t count, size_t offset, enum cudaMemcpyKind kind); cudaError_t cudaMemcpyToSymbolAsync(const char *symbol, const void *src, size_t count, size_t offset, enum cudaMemcpyKind kind, cudaStream_t stream); cudaError_t cudaMemcpyFromSymbolAsync(void *dst, const char *symbol, size_t count, size_t offset, enum cudaMemcpyKind kind, cudaStream_t stream);
Параметр kind задает направление копирования и может принимать только одно из следующих значений:
- cudaMemcpyHostToHost - скопировать из памяти CPU в память CPU;
- cudaMemcpyHostToDevice - скопировать из памяти CPU в память GPU;
- cudaMemcpyDeviceToHost - скопировать из памяти GPU в память CPU;
- cudaMemcpyDeviceToDevice - скопировать из памяти GPU в память GPU.
Параметр stream позволяет организовать несколько потоков команд, если вы не используете эту возможность, то следует использовать поток по умолчанию - 0.
Ниже приводится простейший пример - фрагмент кода, объявляющий массив в константной памяти и инициализирующий этот массив данными из массива, расположенного в памяти CPU.
__constant__ float contsData[256]; float hostData[256]; ... cudaMemcpyToSymbol(constData, hostData,sizeof(data), 0, cudaMemcpyHostToDevice);
Поскольку константная память кэшируется, то она является идеальным местом для размещения небольшого объема часто используемых неизменяемых данных, которые должны быть доступны всем нитям сетки сразу.
На следующем шаге мы рассмотрим пример использования глобальной памяти.
