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

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个线程，我们需要30次4字节的读取才能将数字存储在每个块的共享内存中。

如果没有共享内存，如果每个线程都必须读取这三个变量一次，那么总共读取全局内存的数量将是1024*10*3 = 30720，非常低效。

现在问题来了，我对CUDA还比较新手，不确定是否可以将变量a、b和c的内存转移到每个块的共享内存中，而不必让每个线程从全局内存中读取这些变量并将它们加载到共享内存中，因此最终全局内存读取的总量将是1024*10*3 = 30720，而不是30。

在以下网站上有一个例子：

 __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 次全局内存读取来实现这一点？

- ksm001

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架构）。按照您的建议进行操作没有任何优势。

- talonmies

---

此内容可在CUDA Toolkit v6.5文档的“E.2.5.3.函数参数”章节中找到，属于编程指南。 - JonathanK