传递一个常量共享内存

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__空间。

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

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

注意事项：

1. 有许多 CUDA 示例代码演示了共享内存的更高级用法。

2. 此处的使用不会带来任何好处。共享内存通常用于需要线程间通信的情况，或者用于数据重用的情况。您的代码都没有演示。

3. 还有许多其他方法可以向内核提供常量值，包括常量数组，例如__constant__ 内存。这些中的任何一个是否有益将在很大程度上取决于您的实际用例和访问模式，我认为您所显示的代码并未表示这些。无论如何，这里有很多关于 CUDA 标签的问题，它们讨论了各种持续数据的使用，我相信你可以通过一些搜索找到。

4. __syncthreads() 可以说对于此代码不是必需的。但是在共享内存的许多更典型的用途中它是必要的，所以我选择在这里指出它。在这个特定的代码中，它不是必需的，但是这个特定的代码也不会是共享内存的合理使用。