传递常量共享内存

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 示例代码演示了共享内存的更高级用法，例如6_Performance/transpose.

2. 这里的用法不会带来任何好处。 共享内存通常用于需要线程间通信的情况，或者在存在数据重用的情况下。 您的代码不演示这两种情况。

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

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