Какая разница между общей областью CUDA и глобальной памятью?

Я запутался в том, как использовать общую и глобальную память в CUDA, особенно в отношении следующего:

Когда мы используем cudaMalloc(), получаем ли мы указатель на общий или глобальный память?
Сохраняется ли глобальная память на хосте или устройстве?
Есть ли ограничение размера для одного?
К чему быстрее обращаться?
Сохраняет переменная в общей памяти такая же, как и передача ее адреса через ядро? То есть вместо
```
__global__ void kernel() {
   __shared__ int i;
   foo(i);
}
```
почему не эквивалентно делать
```
__global__ void kernel(int *i_ptr) {
   foo(*i_ptr);
}

int main() {
   int *i_ptr;
   cudaMalloc(&i_ptr, sizeof(int));
   kernel<<<blocks,threads>>>(i_ptr);
}
```

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

Большое спасибо

Ответ 1

Когда мы используем cudaMalloc()

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

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

Глобальная память находится на устройстве. однако существуют способы использования памяти хоста в качестве "глобальной" памяти с использованием сопоставленной памяти, см.: Особенности памяти CUDA Zero Copy, однако это может быть медленной скоростью из-за ограничений скорости передачи по шине.
Есть ли ограничение на размер одного из них?

Размер глобальной памяти зависит от карты к карте, от нуля до 32 ГБ (V100). При этом общая память зависит от вычислительных возможностей. Все, что ниже вычислительной способности 2.x, имеет максимум 16 КБ общей памяти на многопроцессорность (где количество многопроцессоров варьируется от карты к карте). А карты с вычислительной способностью 2x и выше имеют как минимум 48 КБ разделяемой памяти на каждый многопроцессор.

Смотрите https://en.wikipedia.org/wiki/CUDA#Version_features_and_specifications

Если вы используете подключенную память, единственным ограничением является то, сколько хост-компьютера находится в памяти.
Что быстрее?

Что касается необработанных чисел, общая память намного быстрее (общая память ~ 1,7 ТБ/с, а глобальная память ~ XXXGB/с). Однако для того, чтобы что-то сделать, вам нужно чем-то заполнить общую память, вы обычно извлекаете из глобальной памяти. Если доступ к глобальной памяти объединен (неслучайно) и имеет большой размер слова, вы можете достичь скоростей, близких к теоретическому пределу в сотни ГБ/с, в зависимости от карты и ее интерфейса памяти.

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

Кроме того, его часто используют в качестве блокнота для уменьшения нагрузки на регистр, что влияет на количество рабочих групп, которые можно запускать одновременно.
Хранение переменной в разделяемой памяти - это то же самое, что передача ее адреса через ядро?

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

Ответ 2

Содержимое глобальной памяти видны всем потокам сетки. Любой поток может читать и записывать в любое место глобальной памяти.

Общая память является отдельной для каждого блока сетки. Любой поток блока может читать и записывать в общую память этого блока. Поток в одном блоке не может получить доступ к общей памяти другого блока.

cudaMalloc всегда выделяет глобальную память.
Глобальная память находится на устройстве.
Очевидно, что каждая память имеет ограничение по размеру. Глобальная память - это общее количество DRAM используемого вами графического процессора. Например, я использую GTX460M, который имеет 1536 МБ DRAM, поэтому 1536 МБ глобальной памяти. Общая память определяется архитектурой устройства и измеряется для каждого блока. Устройства с вычислительной способностью от 1,0 до 1,3 имеют 16 KB/Block, вычислительные от 2,0 до 7,0 имеют общую память 48 KB/Block.
Совместно используемая память на порядок быстрее, чем глобальная память. Это как локальный кеш, разделяемый между потоками блока.
Нет. Только адреса глобальной памяти могут быть переданы ядру, запущенному с хоста. В первом примере переменная считывается из общей памяти, а во втором - из глобальной памяти.

Обновление:

Устройства Compute Capability 7.0 (архитектура Volta) позволяют выделять общую память до 96 КБ на блок при условии соблюдения следующих условий.

Общая память распределяется динамически
Перед запуском ядра максимальный размер динамической разделяемой памяти задается с помощью функции cudaFuncSetAttribute следующим образом.

__global__ void MyKernel(...)
{
    extern __shared__ float shMem[];
}

int bytes = 98304; //96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes);

MyKernel<<<gridSize, blockSize, bytes>>>(...);

Ответ 3

Общая память CUDA - это память, разделяемая между потоками внутри блока, то есть между блоками в сетке содержимое разделяемой памяти undefined. Его можно считать кэшем L2 с ручным управлением.

Обычно глобальная память хранится на устройстве, но последние версии CUDA (если устройство поддерживает ее) могут отображать память хоста в адресное пространство устройства, вызывая в этом случае передачу DMA на месте с хоста на память устройства.

В зависимости от устройства существует ограничение размера для общей памяти. Он сообщается в возможностях устройства, извлекаемых при перечислении устройств CUDA. Глобальная память ограничена общей памятью, доступной для GPU. Например, GTX680 предлагает 48kiB общей памяти и 2GiB памяти устройства.

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

Сохраняет ли переменная в общей памяти то же самое, что передавать ее адрес через ядро?

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