CUDA共享和全局内存之间有什么区别？

Question

我对如何在CUDA中使用共享和全局内存感到困惑,尤其是对于以下内容:

• 当我们使用时cudaMalloc(),我们得到一个指向共享或全局内存的指针吗？
• 全局内存是否驻留在主机或设备上？
• 是否有任何尺寸限制？
• 哪个访问速度更快？

• 将变量存储在共享内存中与通过内核传递其地址相同吗？也就是说,而不是

  __global__ void kernel() {
     __shared__ int i;
     foo(i);
  }
  

  为什么不等同呢

  __global__ void kernel(int *i_ptr) {
     foo(*i_ptr);
  }
  
  int main() {
     int *i_ptr;
     cudaMalloc(&i_ptr, sizeof(int));
     kernel<<<blocks,threads>>>(i_ptr);
  }
  

关于全局与共享内存中的特定速度问题存在很多问题,但没有一个问题包含何时在实践中使用任何一个内容的概述.

非常感谢

Answer 1

• 当我们使用cudaMalloc()时

  为了将数据存储在可以传回主机的gpu上,我们需要有一个存储的内存,直到它被释放,将全局内存看作具有生命的堆空间直到应用程序关闭或被释放,它是可见的到具有指向该内存区域的指针的任何线程和块.共享内存可以被视为具有生命的堆栈空间,直到内核块完成,可见性仅限于同一块内的线程.所以cudaMalloc用于在全局内存中分配空间.

• 我们得到一个指向共享或全局内存的指针吗？

  您将获得一个指向驻留在全局内存中的内存地址的指针.

• 全局内存是否驻留在主机或设备上？

  全局内存驻留在设备上.但是,有一些方法可以使用映射内存将主机内存用作"全局"内存,请参阅:CUDA归零内存注意事项,但由于总线传输速度限制,可能速度较慢.

• 是否有任何尺寸限制？

  全局内存的大小取决于卡到卡,从无到8GB.共享内存取决于计算能力.低于计算能力2.x的任何东西每个多处理器最多有16KB的共享内存(多处理器的数量因卡而异).计算能力为2.x或更高的卡每个多处理器最多可以有48KB的共享内存.

  如果使用映射内存,唯一的限制是主机在内存中的容量.

• 哪个访问速度更快？

  就原始数字而言,共享内存要快得多(共享内存~1.7TB/s,而全局内存~150GB/s).但是,为了做任何你需要用某些东西填充共享内存的东西,你通常会从全局内存中取出.如果内存访问全局内存是合并的(非随机),您将获得高达150-200GB/s的速度,具体取决于卡及其内存接口.

  共享内存的使用是指您需要在一个线程块内重用已经从全局内存中提取或评估的数据.因此,不要再次从全局内存中提取,而是将其放在同一块中的其他线程的共享内存中以查看和重用.

• 将变量存储在共享内存中与通过内核传递其地址相同吗？

  不,如果您传递任何地址,它始终是全局内存的地址.在主机中,您无法设置共享内存,除非您将其作为常量传递,内核将共享内存设置为该常量,或者将地址传递给全局内存,并在需要时由内核提取.

Answer 2

全局内存的内容对于网格的所有线程都是可见的.任何线程都可以读写全局内存的任何位置.

共享内存对于网格的每个块是独立的.块的任何线程都可以读取和写入该块的共享内存.一个块中的线程无法访问另一个块的共享内存.

1. cudaMalloc 总是分配全局内存.
2. 全局内存驻留在设备上.
3. 显然,每个内存都有一个大小限制.全局内存是您正在使用的GPU的DRAM总量.例如,我使用的是具有1536 MB DRAM的GTX460M,因此全局内存为1536 MB.共享内存由设备体系结构指定,并以每个块为基础进行测量.计算能力1.0到1.3的设备16 KB/Block,计算2.0到7.0具有48 KB/Block共享内存.
4. 共享内存的访问速度比全局内存快.它就像一个块的线程之间共享的本地缓存.
5. 不可以.只有全局内存地址可以传递给从主机启动的内核.在第一个示例中,从共享内存中读取变量,而在第二个示例中,从全局内存中读取变量.

更新:

计算能力设备7.0(Volta Architecture)允许在每个块上分配高达96 KB的共享内存,前提是满足以下条件.

• 共享内存是动态分配的
• 在启动内核之前,使用以下函数指定动态共享内存的最大大小cudaFuncSetAttribute.

__global__ void MyKernel(...)
{
    extern __shared__ float shMem[];
}

int bytes = 98304; //96 KB
cudaFuncSetAttribute(MyKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, bytes);

MyKernel<<<gridSize, blockSize, bytes>>>(...);