传递常量共享内存

Question

我有一个长度为 128 的向量；在整个计算过程中，所有元素都是恒定的。

我喜欢在我的 CUDA 内核中使用这个常数向量。我正在考虑将这个向量存储在共享内存中，并在内核中使用它。我想知道该怎么做？几行代码就好了。

或者这是最好的方法吗？非常感谢。

头顶的我们可以通过全局内存：

__global__ void fun(float* a, float* coeff)
{
  size_t
        i = blockIdx.x * blockDim.x + threadIdx.x;

  if (i >= 128)
       return;

  a[i] *= coeff[i];
}

但这可能不是最好的方法。我想像

__shared__ float coeff[128];

但是如何将 CPU 值复制到此共享内存？我是否将此共享内存传递给我的内核？

Answer 1

__shared__ 内存不能直接从主机代码访问。所以你必须首先通过全局内存将数据传递给它，然后从那里将它复制（使用内核代码）到 __shared__ space.

对您的内核代码进行简单修改以演示该概念如下所示：

__global__ void fun(float* a, float* coeff)
{
  __shared__ float scoeff[128];
  size_t
        i = blockIdx.x * blockDim.x + threadIdx.x;

  if (i >= 128)
       return;
  scoeff[i] = coeff[i];
  __syncthreads();

  a[i] *= scoeff[i];
}

备注：

有许多 CUDA 示例代码演示了共享内存的更高级用法。
这里的用法没有任何好处。共享内存通常用于需要线程间通信的情况，或者用于数据重用的情况。您的代码没有展示任何内容。
有许多其他方法可以为内核提供常量值，包括常量数组，例如 __constant__ 内存。这些是否有益将在很大程度上取决于您的实际用例和访问模式，我认为这些用例和访问模式并未由您显示的代码表示。无论如何，CUDA 标签上有很多问题讨论了各种常量数据的使用，我相信您可以通过一些搜索找到这些问题。
__syncthreads() 可以说对于这段代码来说不是必需的。但在共享内存的许多更典型的用途中它是必要的，所以我选择在这里指出它。在此特定代码中，没有必要，但此特定代码也不是共享内存的明智使用。

传递常量共享内存

Pass a constant shared memory

cuda

vector

shared-memory