CUDA объединила доступ к глобальной памяти

Question

Я прочитал руководство по программированию CUDA, но я упустил одну вещь. Допустим, у меня есть массив 32-битных int в глобальной памяти, и я хочу скопировать его в общую память с объединенным доступом. Глобальный массив имеет индексы от 0 до 1024, и скажем, у меня есть 4 блока по 256 потоков каждый.

__shared__ int sData[256];

Когда осуществляется объединенный доступ?

1.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y];

Адреса в глобальной памяти копируются от 0 до 255, каждый по 32 потока в деформации, так что здесь все в порядке?

2.

sData[threadIdx.x] = gData[threadIdx.x * blockIdx.x+gridDim.x*blockIdx.y + someIndex];

Если someIndex не кратен 32, это не объединяется? Смещенные адреса? Это верно?

Answer 1

То, что вы хотите, в конечном счете, зависит от того, являются ли ваши входные данные одномерным или двумерным массивом, и являются ли ваши сетка и блоки одномерными или двумерными. Самый простой случай - это 1D:

shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + threadIdx.x];

Это слилось. Основное правило, которое я использую, заключается в том, что наиболее быстро меняющаяся координата (threadIdx) добавляется как смещение к смещению блока (blockDim * blockIdx). Конечный результат заключается в том, что шаг индексации между потоками в блоке равен 1. Если шаг увеличивается, то вы теряете объединение.

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

На GT2xx и более ранних GPU все становится сложнее. Но вы можете найти детали этого в руководстве по программированию.

Дополнительные примеры:

Не объединены:

shmem[threadIdx.x] = gmem[blockDim.x + blockIdx.x * threadIdx.x];

Не слились, но не так уж плохо на GT200 и позже:

stride = 2;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

Вообще не слились

stride = 32;
shmem[threadIdx.x] = gmem[blockDim.x * blockIdx.x + stride * threadIdx.x];

Объединенная, 2D сетка, 1D блок:

int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.x] = gmem[blockIdx.y * elementPitch + 
                          blockIdx.x * blockDim.x + threadIdx.x]; 

Объединенная, 2D сетка и блок:

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int elementPitch = blockDim.x * gridDim.x;
shmem[threadIdx.y * blockDim.x + threadIdx.x] = gmem[y * elementPitch + x];

Answer 2

Ваше индексирование в 1 неверно (или намеренно настолько странно, что кажется неправильным), некоторые блоки обращаются к одному и тому же элементу в каждом потоке, поэтому нет возможности объединить доступ в этих блоках.

Доказательство:

Пример:

Grid = dim(2,2,0)

t(blockIdx.x, blockIdx.y)

//complete block reads at 0
t(0,0) -> sData[threadIdx.x] = gData[0];
//complete block reads at 2
t(0,1) -> sData[threadIdx.x] = gData[2];
//definetly coalesced
t(1,0) -> sData[threadIdx.x] = gData[threadIdx.x];
//not coalesced since 2 is no multiple of a half of the warp size = 16
t(1,1) -> sData[threadIdx.x] = gData[threadIdx.x + 2];

Так что это игра "удачи", если блок слился, так что в общем нет

Но правила чтения объединенной памяти не так строги в новых версиях cuda, как раньше.
Но для проблем совместимости вы должны попытаться оптимизировать ядра для самых низких версий cuda, если это возможно.

Вот хороший источник:

http://mc.stanford.edu/cgi-bin/images/0/0a/M02_4.pdf

Answer 3

Правила, по которым доступ может быть объединен, несколько сложны и со временем изменились. Каждая новая архитектура CUDA является более гибкой в ​​том, что она может объединять. Я бы сказал, чтобы не беспокоиться об этом вначале. Вместо этого сделайте доступ к памяти любым удобным для вас способом, а затем посмотрите, что говорит профилировщик CUDA.

Answer 4

Ваши примеры верны, если вы намеревались использовать 1D сетку и геометрию резьбы. Я думаю, что индексирование, которое вы намеревались использовать, [blockIdx.x*blockDim.x + threadIdx.x],

С #1 32 потока в деформации выполняют эту инструкцию "одновременно", так что я считаю, что их запросы, которые являются последовательными и выровненными по 128B (32 x 4), объединяются в архитектурах Tesla и Fermi.

С #2 это немного размыто. Если someIndex равным 1, то он не объединит все 32 запроса в деформации, но может частично слить. Я полагаю, что устройства Fermi объединят доступ для потоков 1-31 в деформации как часть последовательного сегмента памяти 128B (и первые 4B, которые не нужны ни одному потоку, теряются). Я думаю, что устройства с архитектурой Tesla сделали бы такой доступ несвязанным из-за смещения, но я не уверен.

С someIndex как, скажем, 8, у Теслы будут выровненные адреса 32B, и Fermi может сгруппировать их как 32B, 64B и 32B. Но суть в том, в зависимости от значения someIndex и архитектура, что происходит, размыта, и это не обязательно будет ужасно.