在不使用内核的情况下写入 CUDA 中的共享内存答案

【问题标题】：Writing to Shared Memory in CUDA without the use of a kernel在不使用内核的情况下写入 CUDA 中的共享内存
【发布时间】：2023-03-23 02:25:01
【问题描述】：

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

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

对此的任何帮助将不胜感激。提前致谢！

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

编辑：显然这对于共享内存是不可能的。有谁知道只读缓存是否可行？

【问题讨论】：

标签： c++ memory cuda main shared

【解决方案1】：

这是不可能的。填充 shared memory 的唯一方法是在 CUDA 内核中使用线程。

如果您希望内核在启动时可以使用一组（只读）数据，当然可以使用__constant__ memory。可以使用文档中指示的 API 在主机代码上/由主机代码设置此类内存，即cudaMemcpyToSymbol。

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

int myval = constant_data[12];

否则使用普通全局内存，无论是静态的还是动态分配的，使用适当的主机 API 进行初始化（动态：cudaMemcpy，静态：cudaMemcpyToSymbol）。

【讨论】：

那太臭了。您知道是否可以使用计算 3.5 设备中的只读缓存来完成我想要的操作？谢谢。
让我们明确一点，SO 不是聊天服务。如果您有新问题，请将其作为新问题发布。
如果 60% 的线程读取相同的地址，而其他 40% 的线程以相等的比例读取其他 3 个地址，您认为常量内存比全局内存更有效吗？
很难说。确切的访问模式将影响每种情况下的性能。您可以尝试编写一个基准来比较这两种情况；应该不难。

【解决方案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 的所有负载都将完美合并，因此这应该接近最佳效率。

【讨论】：