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

ut6juiuv  于 2023-05-20  发布在  其他
关注(0)|答案(2)|浏览(163)

我想在main()函数中创建一个数组,输入所有正确的值,然后让共享内存中的线程立即使用这个数组。
我在CUDA中查找的每个如何使用共享内存的示例都有线程在向共享数组中写入数据,但我希望在内核启动之前共享数组立即可用。
任何帮助这样做将不胜感激。先谢谢你了!
一些背景:我想要的共享数组永远不会改变,并且所有线程都从它读取。
编辑:显然,这是不可能的共享内存。有谁知道只读缓存是否可能?

webghufk

webghufk1#

不可能的填充共享内存的唯一方法是使用CUDA内核中的线程。
如果您希望内核在启动时可以使用一组(只读)数据,那么当然可以使用__constant__内存。这样的存储器可以使用文档中指示的API在主机代码上/由主机代码设置,即cudaMemcpyToSymbol
__constant__内存实际上只有在每个线程都将在给定的访问周期中访问 * 相同的位置 * 时才有用,例如

int myval = constant_data[12];

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

bf1o4zei

bf1o4zei2#

虽然您请求的特定行为无法自动实现,但这实际上是一个相当常见的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的所有负载将被完美地合并,因此这应该接近最佳效率。

相关问题