cuda学习(1): 内存Memory

0 序言

memory是cuda中很重要的知识点，通常和高性能有关系，你利用的好memory的一些特性就能实现很多高性能的场景。主要对pinned memory、global memory、shared memory进行介绍，其他不常用。

• pinned memory通常指host memory主机内存，global memory 和shared memory属于GPU的内存。

下表是Device内存的类型，常用一般是两种:global memory 和shared memory

1 显卡内存

• 上图是一张显卡的图，这张图不太严谨但能够辅助理解下几个内存之间的关系。中间的绿块是运算单元，用来做计算。Global memory 是运算单元周边的小黑块。Shared memory在运算单元里面有一块区域，所以我们可以简单的认为shared memory它就是片上内存,Global是片外内存。
  
• 上图是电脑主板的图，显卡通过接口插到主板上的PCIE插槽上，内存条插到内存条的插槽上。
• 内存条的内存通常有几种叫法:主机内存、host memory,cpu内存,pinned memory。

1.1 Pinned Memory

对于整个Host Memory内存条而言，操作系统区分为两个大类(逻辑区分，物理上是同一个东西)

• Pageable memory,可分页内存
• Page lock memory,页锁定内存，又叫Pinned Memory

你可以理解为Page lock memory是vip房间，锁定给你一个人用,其他人不可以来占用它。而Pageable memory是普通房间，在酒店不够的时候，选择性的把你的房间腾出来给其他人交换用(内存交换出去，交换到硬盘上，从而之前的内存空间交给其他任务使用，使得性能降低)，这就可以容纳更多人了。造成房间(内存空间)很多的假象，代价是性能降低

总结

• 1. pinned memory具有锁定特性，是稳定不会被交换出去的(这很重要，相当于每次去这个房间找你都一定能找到你)
• 2. pageable memory没有锁定特性，对于第三方设备(比如GPU），去访问时，因为无法感知内存是否被交换，可能得不到正确的数据(每次去房间找，说不准你的房间被人交换了)
• 3. pageable memory的性能比pinned memory差，很可能降低你程序的优先级然后把内存交换给别人用
• 4. pageable memory的策略能使用内存假象，实际8GB但是可以使用15GB，提高程序的运行数量(不是速度), 这种思路就是所谓的虚拟内存
• 5. pinned memory太多，会导致操作系统性能降低(程序运行数量减少)，8GB就只能用8GB。注意不是你的应用程序性能降低，只是你的运行程序减少了
• 6. GPU可以直接访问pinned memory而不能访问pageable memory(因为第二条)，因为pageable memory没有锁定特性，你访问它可能造成错误。

1.2 数据传输到GPU

• Pageable memory 数据传输到GPU,通常要先传给Pinned Memory,然后再传给GPU
• 对于Pinned Memory,如果数据在Pinned Memory上可以直接传输到GPU上，没有中间过程

1.3 显卡访问Pinned Memory轨迹

计算单元要访问Pinned Memory, 通过PICE接口，到主板，再到内存条得到数据，GPU可以直接访问Pinned Memory

原则:

• 1. GPU可以直接访问pinned memory,称之为(DMA Direct Memory Access)
• 2. 对于GPU访问而言，距离计算单元越近，效率越高。最近的是SharedMemory,其次是GlabalMemory,再就是PinnedMemory，所以：
     PinnedMemory < GlobalMemory <SharedMemory
• 3. 代码中，由new,malloc分配的是pageable memory, 由cudaMallocHost分配的才是pinnedMemory,由cudaMalloc分配的是GlobalMemory
• 4. 尽量多用PinnedMemory存储host数据，或者显示处理Host到Device时，用PinnedMemmory做缓存。所以如果你的数据本身在pinnedmemory里面，可以直接传到GPU上，所以尽量多用pinnedmemory存储host数据，但如果你的host数据是在pageablememory上，通过new 或者malloc分配出来的，可以考虑用pinnedmemory做一个二级缓存，先传给pinnedmemory，再传给GPU。这些都是都是提高性能的关键。如果使用pageablememory传递过去，它回去分配一次内存，这种分配释放就很浪费时间。

2.代码

// CUDA 运行时头文件
#include <cuda_runtime.h>

#include <stdio.h>
#include <string.h>

#define checkRuntime(op)  __check_cuda_runtime((op),#op,__FILE__,__LINE__)

bool __check_cuda_runtime(cudaError_t code,const char* op,const char* file,int line)
{
	if(code != cudaSuccess){
		const char* err_name = cudaGetErrorName(code);
		const char* err_message = cudaGetErrorString(code);
		print("runtime_error %s%d %s failed.\n code=%s,message = %s\n",file,line,op,*err_name,*err_message)
		return false;
	}
	return true;
}

int main(){
	int device_id=0;
	checkRuntime(cudaSetDevice(device_id));
	
	// global memory
	float* memory_device=nullptr;
	checkRuntime(cudaMalloc(&memory_device,100*sizeof(float)));//pointer to device
	
	// pageable memory
	float* memory_host=new float[100];
	memory_device[2] = 520.25;
	checkRuntime(cudaMemcpy(memory_device,memory_host,sizeof(float)*100,cudaMemcpyHostToDevice));
	
	//pinned memory 也叫 page lock memory
	float* memory_page_locked=nullptr;
	checkRuntime(cudaMallocHost(&memory_page_locked,100*sizeof(float)));
	checkRuntime(cudaMemcpy(memory_page_locked,memory_device,sizeof(float)*100,cudaMemcpyDeviceToHost));
	
	print("%f\n",memory_page_locked[2]);
	checkRuntime(cudaFreeHost(memory_page_locked));
	delete [] memory_host;
	checkRuntime(cudaFree(memory_device));
}

结果

说明整个内存分配和数据传输都是正确的。