在CUDA中写入共享内存而不使用内核

Question

我想在我的main()函数中创建一个数组，输入所有适当的值，然后让这个数组立即被共享内存中的线程使用。

我已经查找了如何在CUDA中使用共享内存的每个例子都有线程写入共享数组，但我希望我的共享数组在内核启动之前立即可用。

任何帮助做这将是非常感激的。提前感谢!

一些上下文:我想要的共享数组永远不会改变，并且可以被所有线程读取。

编辑:显然这是不可能的共享内存。有人知道只读缓存能不能做到吗?

Answer 1

不可能。填充共享内存的唯一方法是在CUDA内核中使用线程。

如果您希望在启动时内核可以使用一组(只读)数据，那么当然可以使用__constant__内存。这样的内存可以由主机代码使用文档中指出的API来设置，即cudaMemcpyToSymbol。

__constant__内存只有在每个线程在给定的访问周期内访问相同的位置时才真正有用，例如

int myval = constant_data[12];

否则使用普通全局内存，静态或动态分配，使用适当的主机API初始化(dynamic: cudaMemcpy, static: cudaMemcpyToSymbol)。

Answer 2

虽然您所要求的特定行为是不可能自动实现的，但这实际上是一个相当常见的CUDA范例:

首先，让所有线程将表复制到shmem中。

同步线程

访问内核中的数据

如果对数据的访问相当随机，并且希望平均多次访问每个条目，那么这可能是一个很大的性能增益。从本质上讲，您使用shmem作为托管缓存，并将来自DRAM的负载一次聚合到shmem中，以便多次使用。此外，shmem对未合并的负载没有惩罚。

例如，您可以这样编码:

const int buffer_size = 8192; // assume an 8k buffer
float *device_buffer = ; // assume you have a buffer already on the device with the data you want.
my_kernel<<<num_blocks, num_threads, buffer_size>>>(..., buffer_size, device_buffer);
__global__ void my_kernel(..., int buffer_size, const float *device_buffer) {
   extern __shared__ float shmem_buffer[];
   for (int idx = threadIdx.x; idx < buffer_sze; idx += blockDim.x) {
       shmem_buffer[idx] = device_buffer[idx];
   }
   __syncthreads();
   // rest of your kernel goes here.  You can access data in shmem_buffer;
}

换句话说，您只需要显式地对副本进行编码。由于来自DRAM的所有负载将被完美地合并，因此这应该接近于最优效率。