На этом шаге мы рассмотрим разделяемую память.
Следующим типом памяти в CUDA является так называемая разделяемая (shared) память. Разделяемая память, размещенная непосредственно в самом мультипроцессоре и доступная на чтение и запись всем нитям блока, является одним из важнейших отличий CUDA от традиционного GPGPU. Правильное использование разделяемой памяти играет огромную роль в написании эффективных программ для GPU.
Для современных GPU каждый потоковый мультипроцессор содержит 16 Кбайт разделяемой памяти. Она поровну делится между всеми блоками сетки, исполняемыми на мультипроцессоре.
Кроме того, разделяемая память используется также для передачи параметров при запуске ядра на выполнение, поэтому желательно избегать передачи большого объема входных параметров, непосредственно передаваемых ядру в конструкции вызова ядра. При необходимости для передачи большого объема параметров можно воспользоваться кэшируемой константной памятью.
Существует два способа управления выделением разделяемой памяти. Самый простой способ заключается в явном задании размеров массивов, выделяемых в разделяемой памяти (то есть описанных с использованием спецификатора __shared__).
__global__ void incKernal(float *a){ __shared__ float buf[256]; buf[threadIdx.x]=a[blockIdx.x*blockDim.x+threadIdx.x]; }
При этом способе задания компилятор сам произведет выделение необходимого количество разделяемой памяти для каждого блока при запуске ядра.
Кроме того, можно также при запуске ядра задать дополнительный объем разделяемой памяти (в байтах), который необходимо выделить каждому блоку при запуске ядра. Для доступа к такой памяти используется описание массива без явно заданного размера. При запуске ядра начало такого массива будет соответствовать началу дополнительно выделенной разделяемой памяти.
__global__ void incKernal(float *a){ __shared__ float buf[]; buf[threadIdx.x]=a[blockIdx.x*blockDim.x+threadIdx.x]; } incKernal<<<dim3(n/256),dim3(256),k*sizeof(float)>>>(a);
В приведенном выше примере кода каждому блоку будет дополнительно выделено k*sizeof(float) байт разделяемой памяти, которая будет доступна через массив buf. Обратим внимание, что можно задать несколько массивов в разделяемой памяти без явного задания их размера, но тогда в момент выполнения ядра они все будут расположены в начале выделенной блоку дополнительной разделяемой памяти, то есть их начала просто совпадут. В этом случае на программиста ложится ответственность за явное разделение памяти между такими массивами.
Очень часто разделяемая память используется для хранения постоянно используемых значений - вместо того чтобы постоянно извлекать из медленной глобальной памяти, гораздо удобнее их один раз прочесть в разделяемую память.
На следующем шаге мы рассмотрим эффективное использование разделяемой памяти.
