Выделение общей памяти

Question

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

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

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

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

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

Answer 1

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

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

Answer 2

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

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

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

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

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

Answer 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);

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

Answer 4

Из "Руководства по программированию в CUDA C": Конфигурация выполнения указывается путем вставки выражения в форме:

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

где:

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

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

Answer 5

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

__shared__ int a[count];

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

__shared__ int a[100];

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

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

extern __shared__ int a[];

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

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

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