有没有办法为共享内存数组设置默认值？答案

【问题标题】：Is there a way of setting default value for shared memory array?有没有办法为共享内存数组设置默认值？
【发布时间】：2011-06-25 13:42:14
【问题描述】：

考虑以下代码：

__global__ void kernel(int *something) {
extern __shared__ int shared_array[];     

// Some operations on shared_array here.

}

是否可以将整个 shared_array 设置为某个值 - 例如0 - 没有明确解决某个线程中的每个单元格？

【问题讨论】：

标签： cuda

【解决方案1】：

您可以像这样有效地并行初始化共享数组

// if SHARED_SIZE == blockDim.x, eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x) 
    shared_array[i] = INITIAL_VALUE;
__syncthreads();

【讨论】：

当然，只有当你有一个一维块时才会出现这种情况。只是说任何新手都不会落入明显的陷阱。我还想知道 float4 有多少提升，这是在较新的设备上仍然提供的另一个技巧，以及它与这种内存合并类型的 init 相结合提供了多少好处。旁注，如果您在 2D 或 3D 内核中加载，重要的是要知道它们被划分为扭曲，就像 [z][y][x] 的数组在内存中。因此，让 [x] 线程彼此之间写得最接近，而在 [z] 上不同的线程写得最远。
所以我试过了，是的，使用 reinterpret_cast 像devblogs.nvidia.com/parallelforall/… 那样复制 16 字节块仍然会得到更好的结果。此外，对要复制的数据设置 restrict 也很重要 --- 轻松提升 15% 的性能，大约与 16 字节对齐的块一样多。
请注意，内存合并很容易就意味着 16 倍的速度差异。
@ÍhorMé 不知道你为什么说这不适用于 2D 或 3D 块。您需要适当地进行寻址，但无论块尺寸如何，一般方法都有效。
是的，我说的是确切的代码 sn-p。基本的想法是可行的，但是为了用 2D/3D 块来实现它，我们还必须知道块是如何划分成 warp 的，以确保内存能够合并。

【解决方案2】：

没有。共享内存未初始化。您必须以某种方式自己初始化它...

来自 CUDA C 编程指南 3.2，B.2.4.2 节，第 2 段：

__shared__ 变量不能将初始化作为其声明的一部分。

这也会丢弃共享变量的重要默认构造函数。

【讨论】：

【解决方案3】：

是的，你可以。您可以指定块中的第一个线程设置它，而另一个不设置它例如：

extern __shared__ unsigned int local_bin[]; // Size specified in kernel call

if (threadIdx.x == 0) // Wipe on first thread - include " && threadIdx.y == 0" and " && threadIdx.z == 0"  if threadblock has 2 or 3 dimensions instead of 1.
{
    // For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}

// Do stuff unrelated to local_bin here    

__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.

// Do stuff to local_bin here

理想情况下，您应该在调用 syncthreads 之前做尽可能多的工作，因为这允许所有其他线程在 memset 完成之前完成它们的工作 - 显然，这仅在工作有可能具有完全不同的线程时才重要完成时间，例如，如果有条件分支。请注意，对于线程 0“设置”for 循环，您需要将 local_bin 数组的大小作为参数传递给内核，以便知道您正在迭代的数组的大小。

Original concept source

【讨论】：

谢谢。我在最终实现中使用了类似的东西。
这样就失去了并行化的好处，总是尽量使用threadIdx和BlockIdx
我认为您总体上是对的，但这取决于是否可以从中计算出初始化值，或者它是否具有更复杂的数值/等分配模式。