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

Question

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

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

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

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

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

Answer 1

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

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

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

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

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

Answer 2

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

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

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)

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