caffe源码学习 — SyncedMemory模块
SyncedMemory负责异构体系(CPU与GPU合作)下的内存数据交换和同步。CPU负责逻辑性强的事务处理,GPU只负责高度线程化的数据处理。在这个体系中,CPU和GPU之间将会存在大量的数据交换,因此使用该模块进行数据块管理。
1. 锁页内存
1.1 锁页内存和可分页内存
锁页内存(pinned memory, page-locked memory)。驻留在物理内存的存储空间,不与物理硬盘交换,不会被分配到低速的虚拟内存中,能够通过DMA加速与设备端进行通信。
可分页内存(pageable memory),由操作系统API(malloc(), new()
)分配的存储空间,可以交换到物理硬盘。
一方面GPU上的显存都是锁页的(因为GPU上的显存不支持交换到磁盘),GPU无法直接访问硬盘等存储设备,只能与内存进行数据交换;另一方面,操作系统的内存管理模块在进行内存管理时,会对物理内存进行虚拟化,并会使用一部分硬盘空间作为虚拟内存,一部分物理内存中当前未使用的数据会被转移到物理硬盘(需要时再读取回物理内存)。因此,为了保证GPU访问的一直是物理内存,需要对物理内存上的一部分空间进行锁页(page-locked)操作(即该部分数据一直驻留内存,不会被转移到硬盘),这部分空间则称为锁页内存。
锁页内存允许GPU上的DMA控制器在使用主机内存时不用CPU参与。这样数据交换可以与内核执行并行处理,而且减少硬盘存储与物理内存的数据传输。
有以下三种方式提升传输速度:
- 使用分页锁定内存,分页锁定内存和显存之间的拷贝速度大约是6GB/s,普通的分页内存和GPU间的速度大约是3GB/s,(另外:GPU内存间速度是30G,CPU间内存速度是10GB/s),但是这种方法会带来额外的cpu内存间的拷贝时间(CPU需要先把数据从可分页内存拷贝到分页锁定内存)。
- 使用内存映射(Zero Copy)让GPU直接使用CPU的内存,减少主机和设备间内存传输的时间,但是这种方法对于2.2以后的cuda版本未必管用。
- 通过函数
cudaHostRegister()
把普通内存改为分页锁定内存,和①的方法类似,但不会带来额外的cpu内存拷贝时间。
锁页内存比可分页内存有更大的分配及回收开销,但是传输大量数据时,提供更快的传输速度。具体的性能提升还取决于显卡的计算能力。
内存锁页操作对于单GPU的提升可能不显著,但是对于多GPU的并行训练,性能提升是十分显著的,而且能提升多GPU上多模型的稳定性。
但分配过多的page-locked memory, 可能会在内存紧张的情况下, 损害操作系统的整体性能.
1.2. CPU/GPU内存分配与释放
- CPU分配内存
通过C标准库中的malloc
函数完成(释放函数free
)
通过MKL库的mkl_malloc
函数申请内存(释放函数mkl_free
)
调用CUDA中的cudaMallocHost
函数(释放函数cudaFreeHost
) - GPU分配显存
通过CUDA的cudaMalloc
函数(释放函数cudaFree
)
通过CUDA的cudaMallocPitch
函数
通过CUDA的cudaMallocArray
函数
cudaMallocPitch
和cudaMallocArray
都没有cudaMalloc
分配的快,当大于4M时,要慢2-4个数量级。
1.3. CPU和GPU之间的内存传输
在GPU核执行任务前,需要把数据都从CPU内存传输到GPU内存,任务完成后,再将运算结果或处理后的数据传回CPU,常用的函数是cudaMemcpy
。
对于小块数据,差异不大,对于大块数据,Pinned memory提供了2.4倍于Non-Pinned memory的吞吐率。
- 分配的内存被用来从CPU传到GPU,然后从GPU传回CPU,两次传输的size相同;
- 分配的内存全部从CPU传到GPU,然后从GPU只是传递了运算结果回CPU,这种情况下GPU核通常执行的是reduction操作,比如常见的计算平均数。
具体使用哪种方式,要结合数据的大小以及CPU和GPU之间传输的类型,当传输的数据比较大,且均需从CPU传至GPU和从GPU传至CPU时,使用Pinned Memory有比较大的性能。
【参考资料 https://satisfie.github.io/archives/page/2/ 】
1.4 内存申请/释放函数CaffeMallocHost / CaffeFreeHost
在snycedmem.hpp
中定义。
传入一个二级指针,该函数将修改该指针所保存的地址值,使其指向GPU申请的内存。根据是否使用GPU、MKL,分别调用cudaMallocHost / mkl_malloc / malloc
函数。
释放ptr
指向的内存空间。根据是否使用GPU、MKL,分别调用cudaFreeHost / mkl_free / free
函数。
2. 模块逻辑
2.1 模块功能
SyncedMemory
中主要完成数据空间的创建,GPU数据和CPU数据,数据空间的释放,及GPU和CPU数据的同步。
同步是指,当一方的数据发生变化时,另外一方的数据需要更新。比如CPU上的数据发生变化,那么需要同步把GPU上的数据更新。这样就涉及到一个何时更新的问题,更新太频繁,比如一方发生变化,就去更新另外一方,这样会导致额外的开销。
2.2 数据存储状态
SyncedMemory
对象管理着一个数据对象,数据对象是一个tensor. 这个数据对象可能只存在CPU上(主机内存),或者只存在GPU上,或者同时存在两个位置上。
- 只存在CPU上,这时候通过
SyncedMemory
对象的私有指针*cpu_ptr_
就能访问; - 只存在GPU上,这时候通过
SyncedMemory
对象的私有指针*gpu_ptr_
就能访问; - 需要同时存在CPU和GPU上,则数据需要同步。希望的效果是两个不同位置的数据是一致的。
涉及到同步管理的问题,使得在对应设备(cpu/gpu)上读取的数据是最新的。每个SyncedMemory
对象有一个标志数据同步状态的状态机变量head_
.
head_
:是一个SyncedHead
类型的枚举变量,取值为
UNINITIALIZED
, 数据对象的存储空间未被初始化,需要malloc
内存(CPU/GPU上);
HEAD_AT_CPU
, 目前在CPU空间上的数据对象的备份是最新的(GPU上的数据则未必最新);
HEAD_AT_GPU
, 目前在GPU空间上的数据对象的备份是最新的(CPU上的数据则未必最新);
SYNCED
,目前GPU和CPU空间上的数据是同步的,即都是最新的.
2.3 改变存储状态
通过下边图片可以描述存储状态是如何改变的。
【引用自 https://blog.csdn.net/hnshahao/article/details/81218713 】
blob
对象通过调用get方法时cpu_data()、gpu_data()
,就会调用to_cpu()、to_gpu()
函数同步数据,使获得的数据是最新的。
3 SyncedMemory类
接下来具体分析snycedmem.hpp
中定义的SyncedMemory
类。
3.1 同步状态SyncedHead
描述SyncedMemory
对象所管理的这段内存的同步状态,最新数据在CPU,还是GPU,还是CPU和GPU已经同步。定义如下:
3.2 私有成员数据
void* cpu_ptr_; // 位于CPU上的数据的指针
void* gpu_ptr_; // 位于GPU上的数据的指针
size_t size_; // 存储空间的大小
SyncedHead head_; // 同步状态
bool own_cpu_data_; // cpu_ptr_指向的数据空间是否是自己申请的。这里own_cpu_data_和own_gpu_data_不是互斥的关
// 系,可以同时own两个地方的数据这里的own代表的含义是SyncedMemory这个对象是否是真的在管理一段数据空间,有一
// 种可能是SyncedMemory对象所包含的数据指针,是指向的另外一段内存空间,而不是由自己申请的。
bool cpu_malloc_use_cuda_;
bool own_gpu_data_;
int device_;
3.3 get和set方法
1. get方法:cpu_data(),gpu_data(),mutable_cpu_data(),mutable_gpu_data()
const void* cpu_data();
const void* gpu_data();
void* mutable_cpu_data();
void* mutable_gpu_data();
cpu_data(),gpu_data()
返回cpu或者gpu内存指针cpu_ptr_
和gpu_ptr_
(函数内首先会通过to_cpu()
或to_gpu()
把最新数据同步到要获取指针的设备上)。返回值类型为const void*
表示cpu_ptr_
和gpu_ptr_
所指向的内存空间不允许被修改。
mutable_cpu_data()
与cpu_data()
都是返回数据空间的指针,区别在于:前者返回的指针是可变的,即意味着CPU上的数据将是最新的,会设置HEAD_AT_CPU
标志。以mutable_cpu_data()
为例,该函数会设置HEAD_AT_CPU
,并返回一个可被修改内容的指针cpu_ptr_
, 表明目前将可能会在CPU的数据上做改动,即意味着CPU上的数据将是最新的。
2. set方法:set_cpu_data
和set_gpu_data
void set_cpu_data(void* data);
void set_gpu_data(void* data);
以set_cpu_data
为例,若拥有数据(own_gpu_data_=true
),则调用cudaFree
释放cpu_ptr_
指向的数据空间,然后指向data
指针指向的数据,同时own_cpu_data_=flase
,表示该数据不是自己创建的,是共享得来的。
3.4 数据同步和状态改变
主要有to_cpu()、to_gpu()、async_gpu_push
三个函数。
3.4.1 to_cpu()
如果同步状态为“UNINITIALIZED
”,则调用CaffeMallocHost
申请内存,然后调用caffe_memset
进行初始化。
(传入cpu_ptr_
指针的地址;第二个参数为是否使用GPU的标志,两个参数都会在CaffeMallocHost
函数中改变).
如果同步状态为“HEAD_AT_GPU
”, 则调用CaffeMallocHost
申请内存,然后调用caffe_gpu_memcpy
从GPU拷贝数据。
caffe_gpu_memcpy(size_, gpu_ptr_, cpu_ptr_); // cudaMemcpy的封装。判断两个指针是否不相等,是则调用
cudaMemcpy(). //(caffe/src/util/math_functions.cu 中定义)将cpu_ptr_指向的内存中的内容全部设置为0,进行初始化。
3.4.2 to_gpu()
如果同步状态为“`UNINITIALIZED`”,则调用`cudaMalloc`申请显存,并调用`caffe_gpu_memset`来初始化空间数据为0.
CUDA_CHECK(cudaMalloc(&gpu_ptr_, size_));
caffe_gpu_memset(size_, 0, gpu_ptr_); // cudaMemset的封装(caffe/include/util/math_functions.hpp 中定义),
//如果在CPU_ONLY下强行调用,则会调用NO_GPU宏,输出错误信息。
如果同步状态为“HEAD_AT_CPU
”,首先判断gpu_ptr_
是否为空,若是则先调用cudaMalloc
申请显存。
然后调用caffe_gpu_memcpy
拷贝数据到GPU
3.4.3 async_gpu_push
首先确定HEAD_AT_CPU
,然后若gpu_ptr_
为空,则调用cudaMalloc
申请显存。数据异步传输调用cudaMemcpyAsync
。“异步”是指把数据同步到GPU,在同步未完成时就返回,不需要等待完成同步。
#ifndef CPU_ONLY
void async_gpu_push(const cudaStream_t& stream); // GPU模式下。异步传输数据,将数据从cpu拷贝到gpu
#endif