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

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.

在以下网站上有这个例子:

 __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架构)自动给每个块.做你的建议是没有优势的.