Caffe 学习
源码学习
LayerSetUp VS. Reshape
每个Caffe的layer都必须定义一个LayerSetUp()函数,用来执行解析layer parameters以及判断bottom blob shape是否正确等初始化操作。而Reshape()函数主要负责管理blob shape,包括设定top blob的shape以及中间blob buffer的shape等等。但在实际的实现中,两个函数的功能的实现偶有重叠。因此,正确认清两者的subtle difference,才能避免错误的产生并合理组织代码的逻辑。
LayerSetUp()函数在layer的生命周期中只在一开始初始化时被执行一次。而Reshape()函数在每次forward的时候都会被调用,从而动态的分配blob的大小。因此,如果top blob和intermediate blob的形状是固定的,Reshape()的代码块可以合并到LayerSetUp()函数中;但是对于convolution或pooling等对blob shape没有固定要求的layer,blob shape的管理必须在Reshape()函数中实现。That’s why it is called Reshape().
mutable_cpu_data VS. cpu_data
在Caffe的blob类中定义了mutable_cpu_data()
和cpu_data()
两种member function。从它们的定义来看两者的区别其实非常小。
const void* SyncedMemory::cpu_data() {
check_device();
to_cpu();
return (const void*)cpu_ptr_;
}
void* SyncedMemory::mutable_cpu_data() {
check_device();
to_cpu();
head_ = HEAD_AT_CPU;
return (const void*)cpu_ptr_;
}
cpu_data()
返回的是常指针,说明该函数只能用来读取数据。mutable_cpu_data()
函数用来修改数据,并且多了一句head_ = HEAD_AT_CPU;
,用来提醒系统CPU memory中的数据已经被修改过了,需要小心。类似地,mutable_gpu_data()
和gpu_data()
的差别也是如此。
GPU, CUDA and Caffe
GPU的并行架构要从两个方面来理解,一是硬件,二是软件。
硬件
GPU硬件运算模块主要由多个(e.g. 16)个SM (streaming multiprocessor)组成,也叫GPU大核。每个SM由多个(e.g. 192 for Kepler) SP (streaming processor)组成,它是实际执行指令和任务的基本处理单元。因此一个GPU可以包含上千个甚至更多的SP,这正是其强大所在。从这个角度来理解,一个quad-core CPU可以认为是只有一个SM和4个SP的简易版GPU。
软件
在CUDA编程架构里,有thread, block, grip, warp等等概念。
- thread: 执行指令
- block: 数个(512 for CUDA 8.0)threads组成一个block,同一个block中的threads可以通过shared memory通信。
- grid: 数个blocks构成grid。
- warp: GPU执行程序时的调度单位,目前CUDA的warp大小为32,同在一个warp的线程执行相同的指令(即相同的kernel function),这就是所谓的SIMT。
对应关系
程序任务在最小的粒度上由一个个threads组成,一开始CUDA会将threads组成若干个thread blocks(512 threads per block in CUDA 8.0),然后将这些blocks交由SM来运行(8 active blocks per SM in Fermi),同时每个SM上会给这些thread blocks的所有线程分配资源,如shared memory。因为一个SP只可以运行一个thread,而SM上线程多而SP少,所以需要调度来保证公平、避免starvation。因此,在最新的CUDA架构Fermi中,32个threads被分成一个warp(相当于一个排),以warp作为调度和运行的基本单元。在Fermi中一个SM可以同时运行48个active warps,由warp scheduler负责调度(相当于连长调度所有的排轮流执行任务)。因为一个SM的硬件资源在一开始就已经分配给了blocks和threads,每个thread不管是否正在被执行,都在shared_memory上拥有自己的register等等资源,所以上下文切换不会造成任何开销,这也是GPU比CPU更高效的一个因素。 因为一个SM支持的warp数量以及一个warp包含的线程数量都由具体架构指定,所以在Fermi架构中,一个GPU上的resident thread最多只有SM*48*32=SM*1536个。
要在Caffe layer中使用GPU进行并行计算,可以自定义kernel function来封装要并行的指令。函数的申明前要加__device__
或__global__
关键字。调用该function的方法是
kernel_function<Dtype><<<CAFFE_GET_BLOCKS(count), CAFFE_CUDA_NUM_THREADS>>>
,
它会根据可并行单元的数量和每个block最多包含的线程数CAFFE_CUDA_NUM_THREADS
来分恰当数量的block。
CUDA & OpenCL
CUDA和OpenCL的关系都和DirectX与OpenGL的关系很相像。如同DirectX和OpenGL一样,CUDA和OpenCL中,前者是配备完整工具包、针对单一供应商(NVIDIA)的成熟的开发平台,后者是一个开放的标准。虽然两者抱着相同的目标:通用并行计算。但是CUDA仅仅能够在NVIDIA的GPU硬件上运行,而OpenCL的目标是面向任何一种Massively Parallel Processor,期望能够对不同种类的硬件给出一个相同的编程模型。
Reference
http://blog.csdn.net/junparadox/article/details/50540602
https://chrischoy.github.io/research/making-caffe-layer/
http://on-demand.gputechconf.com/gtc-express/2011/presentations/cuda_webinars_WarpsAndOccupancy.pdf
Install Caffe on CentOS
Preliminary
- CentOS 7, x86
- Nvidia Geforce 980
- Cuda 8.0
Problems
cublas_v2.h: No such file or directory. Compilation terminated.
The error is caused by the failure to link to ablas. So install openblas instead:
sudo yum install openblas-devel.x86_64
cmake cannot detect cuda
Need to add cuda include and library path to global path
Compilation
- mkdir build
- cd build
- sudo su
- cmake -DBLAS=open ..
- make -j8
- make install