cuda 学习之内存层次结构

cuda 学习之内存层次结构

要想编写高效的程序，那么一定要对内存结构有比较深刻的认识，就像C/C++里面的堆内存，栈内存，全局存储区，静态存储区，常量区等。Cuda是并行计算框架，而GPU的内存有限，那么如果想编写高效的Cuda程序，首先要对其内存结构有一个简单的认识。

首先我们先上一张图，然后通过解释一些名词和代码来进行解释。

registers：寄存器

它是GPU片上告诉缓存器，执行单元可以以极低的延迟访问寄存器。寄存器的基本单元是寄存器文件（register file），每个寄存器文件大小为32bit。寄存器文件数量虽然客观，但是平均分给并行执行的线程，每个线程拥有的数量就非常有限了。**编程时，不要为每个线程分配过多的私有变量。**下面程序中，aBegin,aEnd,aStep,a等变量都是寄存器变量，每个线程都会维护这些变量。

__global__ void registerDemo(float *B,float *A ,int wA)
{
    int aBegin = wA*BLOCK_SIZE * blockIdx.y;
	int aEnd = aBegin + wA - 1;
	int aStep = BLOCK_SIZE;
	
	for(int a=aBegin;a<=aEnd;a+=aStep)
	{
	    //...
	}
}

local memory：局部存储器

• 位于堆栈中，不在寄存器中的所有内容
• 作用域为特定线程
• 存储在global内存空间中，速度比寄存器慢很多

对于每个线程，局部存储器也是私有的。如果寄存器被消耗完，数据将被存储在局部存储器中。如果每个线程用了过多的寄存器，或声明了大型结构体或数组，或者编译期无法确定数组的大小，线程的私有数据就有可能被分配到local memory中。一个线程的输入和中间变量将被保存在寄存器或者局部存储器中。局部存储器中的数据将被保存在显存中，而不是片上的寄存器或者缓存中，因此对local memory的访问速度比较慢。

shared memory：共享存储器

• SM（SM = streaming multiprocessor）中的内存空间
• 最大48KB
• 作用域是线程块

共享存储器也是GPU片内的告诉存储器。它是一个块可以被同一block中的所有线程访问的可读写存储器。访问共享存储器的速度几乎和访问寄存器一样快。是实现线程间通信的延迟最小的方法。共享存储器可用于实现多种功能，如用于保存共用的计数器（例如计算循环迭代次数）或者block的公共结果（例如规约的结果）。

我们可以动态或者静态的分配shared Memory，其声明即可以在kernel内部也可以作为全局变量。其标识符为：__ shared __。

下面这句话静态的声明了一个2D的浮点型数组：

__shared__ float tile[size_y][size_x];

如果在kernel中声明的话，其作用域就是kernel内，否则是对所有kernel有效。

如果shared Memory的大小在编译器未知的话，可以使用extern关键字修饰，例如下面声明一个未知大小的1D数组：

extern __shared__ int tile[];

由于其大小在编译器未知，我们需要在每个kernel调用时，动态的分配其shared memory，也就是最开始提及的第三个参数：

kernel<<<grid, block, isize * sizeof(int)>>>(...)

应该注意到，只有1D数组才能这样动态使用。

static variable使用shared memory：

#include<iostream>
#include<stdio.h>
 
#if 1
__global__ void example(float *u)
{
	int i = threadIdx.x;
	__shared__ int tmp[4];
	tmp[i] = u[i]; 
	u[i] = tmp[i] * tmp[i] + tmp[3-i] ;
}
#endif
 
#if 1
 
int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u ; 
	size_t size = 4*sizeof(float);
 
	cudaMalloc(&dev_u , size);
	cudaMemcpy(dev_u,host_u,size,cudaMemcpyHostToDevice);
 
	example<<<1,4>>> (dev_u);
 
	cudaMemcpy(host_u , dev_u , size , cudaMemcpyDeviceToHost);
 
	cudaFree(dev_u);
 
	for(int i=0;i<4;i++)
		printf("%f\n",host_u[i]);
	return 0;
}
 
#endif

dynamic variable使用shared memory：

#include<iostream>
#include<stdio.h>
 
__global__ void example(float *u)
{
	int i = threadIdx.x;
  extern __shared__ int tmp[];
	tmp[i] = u[i];
	u[i] = tmp[i] * tmp[i] + tmp[3-i];
}
 
int main()
{
	float host_u[4] = {1,2,3,4};
	float * dev_u;
	size_t size = 4*sizeof(float);
 
	cudaMalloc(&dev_u,size);
	cudaMemcpy(dev_u , host_u ,size , cudaMemcpyHostToDevice);
	example<<<1,4,  size  >>>>(dev_u);
 
	cudaMemcpy(host_u, dev_u,size,cudaMemcpyDeviceToHost);
	cudaFree(dev_u);
	for(int i=0;i<4;i++)
		printf("%f ",host_u[i]);
	return 0;
}

共享内存-Bank Conflict

为了获得较高的内存带宽，共享存储器被划分为多个大小相等的存储器模块，称为bank，可以被同时访问。因此任何跨越b个不同的内存bank的对n个地址进行读取和写入的操作可以被同时进行，这样就大大提高了整体带宽 ——可达到单独一个bank带宽的b倍。但是很多情况下，我们无法充分发挥bank的功能，以致于shared memory的带宽非常的小，这可能是因为我们遇到了bank冲突。

当一个warp中的不同线程访问一个bank中的不同的字地址时，就会发生bank冲突。
如果没有bank冲突的话，共享内存的访存速度将会非常的快，大约比全局内存的访问延迟低100多倍，但是速度没有寄存器快。然而，如果在使用共享内存时发生了bank冲突的话，性能将会降低很多很多。在最坏的情况下，即一个warp中的所有线程访问了相同bank的32个不同字地址的话，那么这32个访问操作将会全部被序列化，大大降低了内存带宽。

NOTE:不同warp中的线程之间不存在什么bank冲突。

共享内存的地址映射方式

要解决bank冲突，首先我们要了解一下共享内存的地址映射方式。
在共享内存中，连续的32-bits字被分配到连续的32个bank中，这就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的。下图中内存地址是按照箭头的方向依次映射的：

上图中数字为bank编号。这样的话，如果你将申请一个共享内存数组(假设是int类型)的话，那么你的每个元素所对应的bank编号就是地址偏移量(也就是数组下标)对32取余所得的结果，比如大小为1024的一维数组myShMem：

• myShMem[4]: 对应的bank id为#4 (相应的行偏移量为0)
• myShMem[31]: 对应的bank id为#31 (相应的行偏移量为0)
• myShMem[50]: 对应的bank id为#18 (相应的行偏移量为1)
• myShMem[128]: 对应的bank id为#0 (相应的行偏移量为4)
• myShMem[178]: 对应的bank id为#18 (相应的行偏移量为5)

下面我介绍几种典型的bank访问的形式。

下面这这种访问方式是典型的线性访问方式(访问步长(stride)为1)，由于每个warp中的线程ID与每个bank的ID一一对应，因此不会产生bank冲突。

下面这种访问虽然是交叉的访问，每个线程并没有与bank一一对应，但每个线程都会对应一个唯一的bank，所以也不会产生bank冲突。

下面这种虽然也是线性的访问bank，但这种访问方式与第一种的区别在于访问的步长(stride)变为2，这就造成了线程0与线程28都访问到了bank 0，线程1与线程29都访问到了bank 2…，于是就造成了2路的bank冲突。我在后面会对以不同的步长(stride)访问bank的情况做进一步讨论。

下面这种访问造成了8路的bank冲突。

这里我们需要注意，下面这两种情况是两种特殊情况：

上图中，所有的线程都访问了同一个bank，貌似产生了32路的bank冲突，但是由于广播(broadcast)机制(当一个warp中的所有线程访问一个bank中的同一个字(word)地址时，就会向所有的线程广播这个字(word))，这种情况并不会发生bank冲突。

同样，这种访问方式也不会产生bank冲突：

这就是所谓的多播机制(multicast)——当一个warp中的几个线程访问同一个bank中的相同字地址时，会将该字广播给这些线程。

NOTE:这里的多播机制(multicast)只适用于计算能力2.0及以上的设备

​ 详细请见 共享内存之bank冲突

constant memory：常数存储器

• 属于全局内存，大小64KB
• 线程请求同一个数据时很快，请求不同的数据时性能下降
• 在运行中不变，所有constant变量的值必须在kernel启动之前从host设置
• __ global __函数参数通过 constant memory穿的到device端， 限定4 KB，即kernel参数通过常量内存传递

它是只读的地址空间。常熟存储器中的数据位于显存，但拥有缓存加速。常数存储器的空间较小，在Cuda程序中用于存储需要频繁访问的只读参数。当来自同一half-warp的线程访问常数存储器中的同一数据时，如果发生缓存命中，那么只需要一个周期就可以获得数据。常数存储器有缓存机制，用以节约带宽，加快访问速度。每个SM拥有8KB的常数存储器缓存。常数存储器是只读的，因此不存在缓存一致性问题。

constant memory的使用：

#include<iostream>
 
using namespace std;
 
__constant__ int devVar = 100;
 
__global__ void xminus(int *a)
{
	int i = threadIdx.x;
	a[i] = devVar+i;
}
 
int main()
{
	int *h_a = (int*)malloc(4*10) ;
	int *d_a ;
	cudaMalloc(&d_a, 4*10) ;
	cudaMemset(d_a, 0, 40) ;
 
	xminus<<<1,4>>>(d_a);
	
	cudaMemcpy(h_a, d_a, 4*10, cudaMemcpyDeviceToHost) ;
 
	for(int i = 0; i < 4 ; i++)
		cout << h_a[i] << " " ;
	cout << endl ;
}

texture memory：纹理存储器

类似constant memory，是只读内存，以某种形式访问的时候可以提升性能。原本是用在OpenGL和DirectX渲染管线中的。
有用的特点：

• 不需考虑要聚合coalescing访问的问题
• 通过“CUDA Array”进行缓存的2D或3D空间的数据位置
• 在1D，2D或3D数组上进行快速插值
• 将整数转换为“unitized”浮点数

global memory：全局存储器

• 独立于GPU核心的硬件RAM
• GPU绝大多数内存空间都是全局内存
• 全局内存的IO是GPU上最慢的IO形式（除了访问host端内存）

全局存储器位于显存（占据了显存的绝大部分），CPU、GPU都可以进行读写访问。整个网格中的任意线程都能读写全局存储器的任意位置由于全局存储器是可写的。全局存储器能够提供很高的带宽，但同时也具有较高的访存延迟。显存中的全局存储器也称为线性内存。线性内存通常使用cudaMalloc()函数分配，cudaFree()函数释放，并由cudaMemcpy()进行主机端与设备端的数据传输。

此外，也可以使用__device__关键字定义的变量分配全局存储器，这个变量应该在所有函数外定义，必须对使用这个变量的host端和device端函数都可见才能成功编译。在定义__device__变量的同时可以对其赋值。

static variable使用global memory：

global_mem_static.cu:

#include<stdio.h>
#include<iostream>
 
__device__ float devU[4];
__device__ float devV[4];
 
//__global__ function 
__global__ void addUV()
{
	int i = threadIdx.x;
	devU[i] += devV[i];
}
 
int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};
	
	int size = 4* sizeof(float);
 
	//cudaMemcpyToSymbol:将数据复制到__constant__或者__device__变量中
	//cudaMemcpyFromSymbol:同上相反
	//cudaMalloc:在设备端分配内存
	//cudaMemcpy:数据拷贝
	//cudaFree():内存释放
	//cudaMemset():内存初始化
	cudaMemcpyToSymbol(devU,hostU,size,0,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(devV,hostV,size,0,cudaMemcpyHostToDevice);
 
	addUV<<<1,4>>>();
 
	cudaMemcpyFromSymbol( hostU,devU,size,0,cudaMemcpyDeviceToHost );
 
	for(int i=0;i<4;i++)
		printf("hostU[%d] = %f\n",i,hostU[i]);
	return 0;
}

dynamic variable使用global memory：

global_mem_dynamic.cu:

#include<iostream>
#include<stdio.h>
 
__global__ void add4f(float *u , float *v)
{
	int i = threadIdx.x;
	u[i] += v[i];
}
 
void print(float * U ,int size)
{
	for(int i=0;i<4;i++)
	{
		printf("U[%d] = %f\n",i,U[i]);
	}
}
 
int main()
{
	float hostU[4] = {1,2,3,4};
	float hostV[4] = {5,6,7,8};
 
	float * devU ;
	float * devV ; 
	int size = sizeof(float) * 4;
	
	//在设备内存上分配空间
	cudaMalloc( &devU,size );
	cudaMalloc( &devV,size );
	//数据拷贝
	cudaMemcpy( devU ,hostU ,size ,cudaMemcpyHostToDevice );
	cudaMemcpy( devV ,hostV ,size ,cudaMemcpyHostToDevice );
 
	add4f<<<1,4>>> (devU,devV);
	//数据返回
	cudaMemcpy(hostU,devU,size,cudaMemcpyDeviceToHost);
 
	print(hostU,size);
	//释放空间
	cudaFree(devV);
	cudaFree(devU);
	return 0;
}

全局内存-对齐与合并

CUDA全局内存的访问是通过”内存事务“实现的，其分类128字节（L1/L2缓存均参与）和32字节（L2缓存参与）两种。

缓存加载（L1+L2）:这种情况下，”内存事务“中加载的缓存粒度是128字节。

1. 对齐合并访问，线程束首地址对齐128字节，且连续访问128字节内存。只需一次”128字节内存事务“即可完成内存请求。效率100%。

   

2. 访问对齐，随机访问128字节内存。由于所请求的内存地址仍然在一个缓存行中，因此，也只需一次”128字节内存事务“即可完成内存请求。效率100%。

3. 非对齐的连续访问，线程束请求的字节未对齐128，而是分布在两个128字节段范围内。由于启用L1缓存，因此加载必须从首地址128的倍数开始，因此需要加载0-127和127-255两个”128字节内存事务“才能完成内存请求。其中一半数据是请求之外的，因此效率=请求加载的全局内存/所需加载的全局内存=50%。 这主要是由于缓存的加载模式导致，因为缓存不是一次性只加载一个数据，而是一批数据，如128字节，而且又必须保证对齐操作，这就导致了加载的浪费。

4. 同一warp中线程只访问一个地址。只需要一个”128内存事务“就可以完成请求，但是效率却非常低。因为所需4字节，而加载了128字节。效率为4/128=3.125%。

5. 最坏的情况，warp中32线程所请求的内存全部分散，因此加载的”内存事务“的可能在0-32之间不等。完成一次内存请求最差需要进行32次”内存事务“，而加载的128字节中，却只有4字节是warp所需的。

非缓存加载（L2）:这种情况下，”内存事务“中加载的缓存粒度是32字节，这是比128字节更细粒度的加载，会对非对齐或非合并的访问带来好处。类似地，对比上述情况，逐一分析。

1. 对齐合并访问，一个warp的128字节的请求，需要4个32字节”内存事务“完成。效率100%。

2. 对齐，但访问是不连续。所需地址将占用4个内存段，不会产生加载浪费。效率100%。

3. 非对齐的连续访问。由于请求内存首地址没有对齐128字节，请求的地址最多落在5个内存段中，总线利用率至少为80%，相比缓存加载50%有了很大的提供，这主要是由于加载了更少的未请求字节。

4. warp中线程束访问一个内存地址。效率为4/32=12.5。非缓存加载要优于缓存加载。

   

5. warp中线程的请求内存地址全部分散，则需要的”内存事务“的也是0-32个，但每个内存事务是32字节，而不再是128字节。这也是非缓存加载优于缓存加载的地方。

host memory :主机端内存

主机端内存，即CPU对应的我们普通意义上的内存。主机端内存分为两种：可分页内存（pageable memory）和页锁定（page-locked 或pinned）内存。可分页内存即为通过操作系统API（malloc()，new()）分配的存储器空间：而页锁定内存始终不会被分配到低俗的虚拟内存中，能够保证存在于屋里内存中，并且能够通过DMA加速与设备端的通信。

总结

下面一张图可以总结这些内存与线程的关系

参考文献：
https://blog.csdn.net/qq_17239003/article/details/79038333
https://blog.csdn.net/lavorange/article/details/20465869
https://segmentfault.com/a/1190000007533157?utm_source=tag-newest