淘先锋技术网

首页 1 2 3 4 5 6 7

    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函数
    cudaMallocPitchcudaMallocArray都没有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中定义。

传入一个二级指针,该函数将修改该指针所保存的地址值,使其指向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上,或者同时存在两个位置上。

  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已经同步。定义如下:

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_dataset_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