CUDA：将相同的内存位置流式传输到所有线程

Question

这是我的问题：我有很多双打（一组77.500个双打）存储在cuda中。现在，我需要大量线程来对该数组进行一系列操作。每个线程将必须读取该数组的SAME元素，执行任务，将结果存储在共享内存中以及读取该数组的下一个元素。请注意，每个线程将必须同时从同一内存位置读取（仅读取）。所以我想知道：有什么方法可以只读取一次内存就向所有线程广播相同的double吗？读很多次是没有用的...任何想法吗？

Answer 1

这是常见的优化。想法是使每个线程与其blockmate合作以读取数据：

// 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（或者更多的整数因子）以消耗整个对象。当您想以同步方式将数据存储回时，可以应用相同的技术。