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参与。这样数据交换可以与内核执行并行处理，而且减少硬盘存储与物理内存的数据传输。
    有以下三种方式提升传输速度：

1. 使用分页锁定内存，分页锁定内存和显存之间的拷贝速度大约是6GB/s，普通的分页内存和GPU间的速度大约是3GB/s，（另外：GPU内存间速度是30G,CPU间内存速度是10GB/s），但是这种方法会带来额外的cpu内存间的拷贝时间（CPU需要先把数据从可分页内存拷贝到分页锁定内存）。
2. 使用内存映射（Zero Copy）让GPU直接使用CPU的内存，减少主机和设备间内存传输的时间，但是这种方法对于2.2以后的cuda版本未必管用。
3. 通过函数cudaHostRegister()把普通内存改为分页锁定内存，和①的方法类似，但不会带来额外的cpu内存拷贝时间。

    锁页内存比可分页内存有更大的分配及回收开销，但是传输大量数据时，提供更快的传输速度。具体的性能提升还取决于显卡的计算能力。
    内存锁页操作对于单GPU的提升可能不显著，但是对于多GPU的并行训练，性能提升是十分显著的，而且能提升多GPU上多模型的稳定性。
    但分配过多的page-locked memory, 可能会在内存紧张的情况下, 损害操作系统的整体性能.

1.2. CPU/GPU内存分配与释放

1. CPU分配内存
   通过C标准库中的malloc函数完成（释放函数free）
   通过MKL库的mkl_malloc函数申请内存（释放函数mkl_free）
   调用CUDA中的cudaMallocHost函数（释放函数cudaFreeHost）
2. 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的吞吐率。

1. 分配的内存被用来从CPU传到GPU，然后从GPU传回CPU，两次传输的size相同；
2. 分配的内存全部从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中定义。

inline void CaffeMallocHost(void** ptr, size_t size, bool* use_cuda)

传入一个二级指针，该函数将修改该指针所保存的地址值，使其指向GPU申请的内存。根据是否使用GPU、MKL，分别调用cudaMallocHost / mkl_malloc / malloc 函数。

inline void CaffeFreeHost(void* ptr, bool use_cuda)

释放ptr指向的内存空间。根据是否使用GPU、MKL，分别调用cudaFreeHost / mkl_free / free 函数。

2. 模块逻辑

2.1 模块功能

    SyncedMemory中主要完成数据空间的创建，GPU数据和CPU数据，数据空间的释放，及GPU和CPU数据的同步。
    同步是指，当一方的数据发生变化时，另外一方的数据需要更新。比如CPU上的数据发生变化，那么需要同步把GPU上的数据更新。这样就涉及到一个何时更新的问题，更新太频繁，比如一方发生变化，就去更新另外一方，这样会导致额外的开销。

2.2 数据存储状态

    SyncedMemory对象管理着一个数据对象，数据对象是一个tensor. 这个数据对象可能只存在CPU上（主机内存），或者只存在GPU上，或者同时存在两个位置上。

1. 只存在CPU上，这时候通过SyncedMemory对象的私有指针*cpu_ptr_就能访问;
2. 只存在GPU上，这时候通过SyncedMemory对象的私有指针*gpu_ptr_就能访问;
3. 需要同时存在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已经同步。定义如下：

enum SyncedHead { UNINITIALIZED, HEAD_AT_CPU, HEAD_AT_GPU, SYNCED };

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()

void to_cpu(); // 同步至CPU，使CPU上数据最新

如果同步状态为“UNINITIALIZED”，则调用CaffeMallocHost申请内存，然后调用caffe_memset进行初始化。

CaffeMallocHost(&cpu_ptr_, size_, &cpu_malloc_use_cuda_);

（传入cpu_ptr_指针的地址；第二个参数为是否使用GPU的标志，两个参数都会在CaffeMallocHost函数中改变）.

caffe_memset(size_, 0, cpu_ptr_); // memset的封装（caffe/include/util/math_functions.hpp 中定义）

如果同步状态为“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()

void to_gpu(); // 同步至GPU，使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