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

Question

请给我一个非常简单的使用共享内存的例子，涉及到IT技术。CUDA C编程指南中提供的示例包含了许多无关的细节。

例如，如果我将一个大数组复制到设备全局内存中，并希望对每个元素进行平方运算，那么如何利用共享内存来加速呢？或者在这种情况下不使用共享内存是否更好？

- Tudor

Answer 1

在你提到的特定情况下，共享内存是没有用处的，原因如下：每个数据元素只使用一次。要使共享内存有用，必须多次使用传输到共享内存中的数据，并使用良好的访问模式，才能发挥其作用。这很简单：仅从全局内存读取需要1次全局内存读取和0次共享内存读取；先将其读入共享内存需要1次全局内存读取和1次共享内存读取，这需要更长的时间。 这里有一个简单的例子，其中块中的每个线程计算相应值的平方，加上其左右邻居的平均值的平方：

  __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;
  }

请注意，这被设计为仅使用一个块运行。扩展到更多块应该很简单。假设块维度为（1024，1，1），网格维度为（1，1，1）。

- Patrick87

---

在开始计算共享内存中的邻居数据之前，您是否忘记了同步屏障？ - pQB

我最喜欢用来解释共享内存的形象是一次F1赛车停车换胎。操作在一个相对受限制的环境中进行，让一群人可以并行地协作完成同一辆赛车的维修和保养任务，从而让任务得以更快地完成。视频链接：www.youtube.com/watch?v=UUvagsM176o - ArchaeaSoftware

非常感谢这个非常有帮助的例子，第一行涉及tmp的末尾多了一个括号，所以我进行了编辑。 - John Powell

@ArchaeaSoftware 这个例子并不相关，我想说！ - undefined

Answer 2

将共享内存看作是一种“显式管理的缓存”-只有在需要多次访问数据时才有用，可以在同一个线程内或者在同一个“块”内的不同线程中访问。如果你只需要访问数据一次，那么共享内存对你没有帮助。

- Paul R

---

例如，在数组平方问题中，缓存不起作用，但如果我在矩阵乘法中缓存A和B，它会起作用，因为它们被多次重复使用。 - Tudor

@Tudor：确切地说——任何在常规应用中可以从L1缓存中获益的东西，都有可能在CUDA应用中通过使用共享内存得到好处。所以，如果你只是读取一个值，将其平方然后写出，那么两种情况下都没有好处。 - Paul R