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

Question

这是我的问题:我有相当大的 double 集(它是 77.500 个 double 的数组)要存储在 cuda 中的某个位置。现在，我需要一大堆线程来顺序地对该数组执行一系列操作。每个线程都必须读取该数组的相同元素，执行任务，将结果存储在共享内存中并读取数组的下一个元素。请注意，每个线程都必须同时从同一内存位置读取(仅读取)。所以我想知道:有没有什么方法可以通过一次内存读取来将相同的 double 广播到所有线程？读很多遍是没有用的......有什么想法吗？

Answer 1

这是一个常见的优化。这个想法是让每个线程与其 block 伙伴合作读取数据:

// 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(或者可能更多的整数因子)元素以消耗整个元素。当您想以同步方式存储数据时，可以使用相同的技术。