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

Я прочитал руководство по программированию CUDA, но пропустил одну вещь. Допустим, у меня есть массив 32-битных целых чисел в глобальной памяти, и я хочу скопировать его в общую память с объединенным доступом. Глобальный массив имеет индексы от 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, он не объединяется? Неверные адреса? Это правильно?