c - CUDA的__shared__内存什么时候有用？

Question

有人可以帮我举一个关于如何使用共享内存的非常简单的例子吗？ Cuda C 编程指南中包含的示例似乎被不相关的细节弄得一团糟。

例如，如果我将一个大数组复制到设备全局内存并想对每个元素进行平方，如何使用共享内存来加快速度？还是在这种情况下没有用？

Answer 1

在您提到的特定情况下，共享内存没有用，原因如下:每个数据元素仅使用一次。要使共享内存有用，您必须多次使用传输到共享内存的数据，使用良好的访问模式，才能发挥作用。原因很简单:仅从全局内存读取需要 1 次全局内存读取和零次共享内存读取；首先将其读入共享内存需要 1 次全局内存读取和 1 次共享内存读取，这需要更长的时间。

这是一个简单的例子，其中 block 中的每个线程计算相应的值，平方，加上它的左右邻居的平均值，平方:

  __global__ void compute_it(float *data)
  {
     int tid = threadIdx.x;
     __shared__ float myblock[1024];
     float tmp;

     // load the thread's data element into shared memory
     myblock[tid] = data[tid];

     // ensure that all threads have loaded their values into
     // shared memory; otherwise, one thread might be computing
     // on unitialized data.
     __syncthreads();

     // compute the average of this thread's left and right neighbors
     tmp = (myblock[tid > 0 ? tid - 1 : 1023] + myblock[tid < 1023 ? tid + 1 : 0]) * 0.5f;
     // square the previousr result and add my value, squared
     tmp = tmp*tmp + myblock[tid] * myblock[tid];

     // write the result back to global memory
     data[tid] = tmp;
  }

请注意，这预计仅使用一个 block 即可工作。扩展到更多 block 应该很简单。假设 block 维度 (1024, 1, 1) 和网格维度 (1, 1, 1)。