从零开始山寨Caffe·贰：主存模型

你左手是内存，右手是显存，内存可以打死显存，显存也可以打死内存。

　　　　　　　　　　　　　　　　　　　　　　 　　　　　　—— 请协调好你的主存

从硬件说起

物理之觞

大部分Caffe源码解读都喜欢跳过这部分，我不知道他们是什么心态，因为这恰恰是最重要的一部分。

内存的管理不擅，不仅会导致程序的立即崩溃，还会导致内存的泄露，当然，这只针对传统CPU程序而言。

由于GPU的引入，我们需要同时操纵俩种不同的存储体：

一个受北桥控制，与CPU之间架起地址总线、控制总线、数据总线。

一个受南桥控制，与CPU之间仅仅是一条可怜的PCI总线。

一个传统的C++程序，在操作系统中，会被装载至内存空间上。

一个有趣的问题，你觉得CPU能够访问显存空间嘛？你觉得你的默认C++代码能访问显存空间嘛？

结果显然是否定的，问题就在于CPU和GPU之间只存在一条数据总线。

没有地址总线和控制总线，你除了让CPU发送数据拷贝指令外，别无其它用处。

这不是NVIDIA解决不了，AMD就能解决的问题。除非计算机体系结构再一次迎来变革，

AMD和NVIDIA的工程师联名要求在CPU和GPU之间追加复杂的通信总线用于异构程序设计。

当然，你基本是想多了。

环境之艰

可怜的数据总线，加大了异构程序设计的难度。

于是我们看到，GPU的很大一部分时钟周期，用在了和CPU互相交换数据。

也就是所谓的“内存与显存之间友好♂关系”。

你不得不接受一个事实：

GPU最慢的存储体，也就是片外显存，得益于镁光的GDDR技术，目前家用游戏显卡的访存速度也有150GB/S。

而我们可怜的内存呢，你以为配上Skylake后，DDR4已经很了不起了，实际上它只有可怜的48GB/S。

那么问题来了，内存如何去弥补与显存的之间带宽的差距？

答案很简单：分时、异步、多线程。

换言之，如果GPU需要在接下来1秒内，获得CPU的150GB数据，那么CPU显然不能提前一秒去复制。

它需要提前3秒、甚至4秒。如果它当前还有其它串行任务，你就不得不设个线程去完成它。

这就是新版Caffe增加的新功能之一：多重预缓冲。

设置于DataLayer的分支线程，在GPU计算，CPU空闲期间，为显存预先缓冲3~4个Batch的数据量，

来解决内存显存带宽不一致，导致的GPU时钟周期浪费问题，也增加了CPU的利用率。

最终，你还是需要牢记一点：

不要尝试以默认的C++代码去访问显存空间，除非你把它们复制回内存空间上。

否则，就是一个毫无提示的程序崩溃问题（准确来说，是被CPU硬件中断了【微机原理或是计算机组成原理说法】)

编程之繁

在传统的CUDA程序设计里，我们往往经历这样一个步骤：

->计算前

cudaMalloc(....) 【分配显存空间】

cudaMemset(....)　　 【显存空间置0】

cudaMemcpy(....)　 【将数据从内存复制到显存】

->计算后

cudaMemcpy(....) 【将数据从显存复制回内存】

这些步骤相当得繁琐，你不仅需要反复敲打，而且如果忘记其中一步，就是毁灭性的灾难。

这还仅仅是GPU程序设计，如果考虑到CPU/GPU异构设计，那么就更麻烦了。

于是，聪明的人类就发明了主存管理自动机，按照按照一定逻辑设计状态转移代码。

这是Caffe非常重要的部分，称之为SyncedMemory(同步存储体)。

主存模型

状态转移自动机

#ifndef CPU_ONLY
void SyncedMemory::async_gpu_data(const cudaStream_t& stream){
CHECK(head_ == HEAD_AT_CPU);
//    first allocating memory
if (gpu_ptr == NULL){
dragonGpuMalloc(&gpu_ptr, size_);
own_gpu_data = true;
}
const cudaMemcpyKind kind = cudaMemcpyHostToDevice;
CUDA_CHECK(cudaMemcpyAsync(gpu_ptr, cpu_ptr, size_, kind, stream));
head_ = SYNCED;
}
#endif

★void SyncedMemory::async_gpu_data()
异步流的底层代码接受一个异步流作为参数，使用cudaMemcpyAsync()向GPU提交复制任务。

它等效于HEAD_AT_CPU+to_gpu()，所以需要更新同步标记。

完整代码将在DataLayer中完成，该函数将由多线程调用。

完整代码

synced_mem.hpp:

https://github.com/neopenx/Dragon/blob/master/Dragon/include/synced_mem.hpp

synced_mem.cpp:

https://github.com/neopenx/Dragon/blob/master/Dragon/src/synced_mem.cpp