Это обычная оптимизация. Идея состоит в том, чтобы заставить каждый поток взаимодействовать с его блочными товарищами для чтения данных:
// choose some reasonable block size
const unsigned int block_size = 256;
__global__ void kernel(double *ptr)
{
__shared__ double window[block_size];
// cooperate with my block to load block_size elements
window[threadIdx.x] = ptr[threadIdx.x];
// wait until the window is full
__syncthreads();
// operate on the data
...
}
Вы можете итеративно "скользить" окном по массиву block_size
(или, может быть, с некоторым целым числом больше) элементов за раз, чтобы поглотить все это. Тот же метод применяется, когда вы хотите хранить данные обратно синхронизированным способом.