CUDA精讲（3）-- GPU体系结构之内存层次结构

内存层次结构

内存层次结构是GPU当中非常重要的一部分，对于性能的优化至关重要，内存的层次结构可以分成一下三个部分来讲：

• Register & Local memory
• Shared memory
• Global & constant memory

1. GPU内存层次概览
   

• 每个thread的local variable对应：register、local memory
• 每个block中的shared variable对应一个shared memory
• 每个grid中的globel varible对应Global memory、constant memory
  

2. register特性:

• register 存在bank冲突
• register不能被索引，也就是说在register中定义一个数组的话会被分配到local memory当中，而local memory其实就是global memroy
• 当register被占满的时候，会使用local memory

3. local memory
   物理上就是global memory，不同于global memory的地方在于：

• 在编译阶段就确定了memory的地址
• local memory可以被cache，L1 cache存在share memory中，L2 cache 存在global memory中。
• 对于local variable定义的array，必须使用静态分配例如int a[10]

4. shared memory特性：

• 几乎和register一样快
• 所有block中的线程都可以通过shared memory进行通信
• 不同block中的线程，只允许通过global memory进行通信
• shared memory只能是一维的，shared varible的动态定义为：

extern __shared__ float A[];

5. global & constant memory
   通过cudaMalloc分配的内存存在于global内存中，如下图：
   

6. constant memory：
   是一个存放只读的数据，通过cudaMemcpyToSymbol & cudaMemcpyFromSymbol两个函数来移动数据。每一个Sm拥有自己的一个大小为64k的constant memory，并且还存在一个8k的cache