Выделение разделяемой памяти

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

__global__ void Kernel(const int count)
{
    __shared__ int a[count];
}

и я получаю сообщение об ошибке

ошибка: выражение должно иметь постоянное значение

количество постоянное Почему я получаю эту ошибку? И как я могу обойти это?

Ответ 1

const не означает "постоянный", это означает "только для чтения".

Постоянное выражение - это то, значение которого известно компилятору во время компиляции.

Ответ 2

CUDA поддерживает динамическое распределение разделяемой памяти. Если вы определите ядро так:

__global__ void Kernel(const int count)
{
    extern __shared__ int a[];
}

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

Kernel<<< gridDim, blockDim, a_size >>>(count)

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

Ответ 3

вариант один: объявить общую память с постоянным значением (не то же самое, что const)

__global__ void Kernel(int count_a, int count_b)
{
    __shared__ int a[100];
    __shared__ int b[4];
}

вариант два: динамически объявлять общую память в конфигурации запуска ядра:

__global__ void Kernel(int count_a, int count_b)
{
    extern __shared__ int *shared;
    int *a = &shared[0]; //a is manually set at the beginning of shared
    int *b = &shared[count_a]; //b is manually set at the end of a
}

sharedMemory = count_a*size(int) + size_b*size(int);
Kernel <<<numBlocks, threadsPerBlock, sharedMemory>>> (count_a, count_b);

Примечание. Указатели на динамически разделяемую память все указаны с тем же адресом. Я использую два общих массива памяти, чтобы проиллюстрировать, как вручную настроить два массива в общей памяти.

Ответ 4

Из "Руководства по программированию CUDA C":

Конфигурация выполнения указывается путем вставки выражения в форме:

<<<Dg, Db, Ns, S>>>

где:

Dg имеет тип dim3 и определяет размер и размер сетки...
Db имеет тип dim3 и определяет размер и размер каждого блока...
Ns имеет тип size_t и определяет количество байтов в разделяемой памяти, которое динамически выделяется на блок для этого вызова в дополнение к статически выделенной памяти. Эта динамически распределенная память используется любой из переменных, объявленных как внешний массив, как упомянуто в __shared__; Ns - необязательный аргумент, который по умолчанию равен 0;
S имеет тип cudaStream_t и определяет связанный поток...

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

Ответ 5

Вы не можете объявить общую переменную, как это.

__shared__ int a[count];

хотя, если вы достаточно уверены в максимальном размере массива a, тогда вы можете прямо объявить, как

__shared__ int a[100];

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

Есть хорошее решение этой проблемы, чтобы объявить

extern __shared__ int a[];

и выделения памяти при вызове ядра из памяти, например

Kernel<<< gridDim, blockDim, a_size >>>(count)

но вам также следует беспокоиться, потому что, если вы используете больше памяти в блоках, чем вы назначаете в ядре, вы получите неожиданные результаты.