cuda - 通过共享内存传递内核参数是否值得？

Question

假设我们有一个数组 int * data ，每个线程将访问该数组的一个元素。由于此数组将在所有线程之间共享，因此将保存在全局内存中。

让我们创建一个测试内核:

 __global__ void test(int *data, int a, int b, int c){ ... }

我肯定知道 data数组将在全局内存中，因为我使用 cudaMalloc 为该数组分配了内存.现在至于其他变量，我看到了一些示例，它们在不分配内存的情况下立即将整数传递给内核函数。在我的例子中，这样的变量是 a b和 c .

如果我没记错的话，即使我们没有直接打电话 cudaMalloc为每三个整数分配4个字节，CUDA会自动为我们做，所以最后变量 a b和 c将分配在全局内存中。

现在这些变量只是辅助，线程只读取它们而没有别的。

我的问题是，将这些变量转移到共享内存不是更好吗？

我想如果我们有例如 10块 1024线程，我们需要 10*3 = 30阅读 4字节，以便将数字存储在每个块的共享内存中。

在没有共享内存的情况下，如果每个线程必须读取所有这三个变量一次，则全局内存读取总量将为 1024*10*3 = 30720这是非常低效的。

现在问题来了，我对 CUDA 有点陌生，我不确定是否可以为变量转移内存 a b和 c到每个块的共享内存，而不需要每个线程从全局内存中读取这些变量并将它们加载到共享内存中，因此最终全局内存读取的总量将为 1024*10*3 = 30720而不是 10*3 = 30 .

上以下 website有这个例子:

 __global__ void staticReverse(int *d, int n)
 {
    __shared__ int s[64];
    int t = threadIdx.x;
    int tr = n-t-1;
    s[t] = d[t];
    __syncthreads();
   d[t] = s[tr];
 }

这里每个线程在共享变量 s 中加载不同的数据.所以每个线程，根据它的索引，在共享内存中加载指定的数据。

就我而言，我只想加载变量 a b和 c到共享内存。这些变量总是相同的，它们不会改变，所以它们与线程本身没有任何关系，它们是辅助的，被每个线程用来运行一些算法。

我应该如何解决这个问题？是否可以仅通过执行 total_amount_of_blocks*3 来实现此目的？全局内存读取？

Answer 1

GPU 运行时已经以最佳方式执行此操作，而您无需执行任何操作(并且您对参数传递在 CUDA 中的工作方式的假设是不正确的)。这是目前发生的情况:

在计算能力 1.0/1.1/1.2/1.3 设备中，内核参数由共享内存中的运行时传递。

在计算能力 2.x/3.x/4.x/5.x/6.x 设备中，内核参数由运行时在保留的常量内存组中传递(它有一个带有广播的专用缓存)。

所以在你假设的内核中

__global__ void test(int *data, int a, int b, int c){ ... }

data , a , b , 和 c都自动按值传递给共享内存或常量内存(取决于 GPU 架构)中的每个块。做你提议的事情没有任何好处。