[资料整理]Caffe:GPU Optimization简介

CUDA(Compute Unified Device Architecture)，是显卡厂商NVIDIA推出的运算平台。 CUDA™是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构ISA以及GPU内部的并行计算引擎。 开发人员现在可以使用C语言来为CUDA™架构编写程序，所编写出的程序于是就可以在支持CUDA™的处理器上以超高性能运行。下面对简要介绍CUDA以及如何在caffe中使用gpu计算。

CUDA简介

Programming model

1.基于CUDA的kernel举例：

Based on  NVIDIA_SAMPLES/SIMPLE/.vecAdd.cu// Kernel definition__global__ void   VecAdd(float* A, float* B, float* C) {  int i = threadIdx.x;  C[i] = A[i] + B[i];}int   main() {   ...   // Kernel invocation with N threads  int numBlocks = 1;  int threadsPerBlock = N;  VecAdd<<< numBlocks, threadsPerBlock >>>(A, B, C);  ...

kernel与普通函数的区别是：

１）在函数定义前加上_global_

２）调用的时候指定blocks的数量和每个blocks中thread的数量

2.Programming model:Grid/Block/Thread

kernel:当执行一个kernel时，相当于在执行一个grid中的thread blocks。

thread block:一个thread block由多个thread组成。

在一个thread block中的threads彼此协作：

(1)在一个thread block中的threads的执行是同步的。

(2)通过低延迟共享内存来高效地共享数据。

注意：

(1)所有的thread共享全局内存。

(2)不同的thread block中的thread不能彼此协作。

Threads和blocks都有ID，例如：Block ID: 1D or 2D；　Thread ID：1D，2D or 3D.

其中，每个blocks中的threads的最大值为1024.

调用kernels的时候要指定blocks的数量和每个blocks中thread的数量：

__global__ void KernelFunc(...);dim3   DimGrid(100, 50);    // 5000 thread blocks dim3   DimBlock(4, 8, 8);   // 256 threads per block KernelFunc<<<DimGrid,DimBlock>>>(...);

3.CUDA SDK

4.CUDA Makefile example

GCC := g++NVCC := nvcc -ccbin $(GCC)CCFLAGS     := -gNVCCFLAGS   := -m64 -g -GLDFLAGS     :=ALL_CCFLAGS := $(NVCCFLAGS) $(addprefix -Xcompiler ,$(CCFLAGS))ALL_LDFLAGS := $(ALL_CCFLAGS) $(addprefix -Xlinker ,$(LDFLAGS))# Common includes and paths for CUDAINCLUDES  := -I../common/incLIBRARIES :=# CUDA code generation flagsGENCODE_SM30    := -gencode arch=compute_30,code=sm_30GENCODE_SM50    := -gencode arch=compute_50,code=sm_50GENCODE_FLAGS   :=  $(GENCODE_SM30)  $(GENCODE_SM50)# Target rulesall:  vectorAddvectorAdd.o:vectorAdd.cu$(NVCC) $(INCLUDES) $(ALL_CCFLAGS) $(GENCODE_FLAGS) -o $@ -c $<vectorAdd: vectorAdd.o$(NVCC) $(ALL_LDFLAGS) $(GENCODE_FLAGS) -o $@ $+ $(LIBRARIES)run: all./vectorAddclean:rm -f vectorAdd vectorAdd.o

5.应用举例

1)矩阵相加：只利用一个block

// Kernel definition__global__  void  MatAdd(floatA[N][N], floatB[N][N],  float  C[N][N]) {    int  i = threadIdx.x;    int  j = threadIdx.y;    C[i][j] = A[i][j] + B[i][j];}int  main() {   …  // Kernel invocation with one block of N * N * 1 threads  int     numBlocks = 1;  dim3 threadsPerBlock(N, N);  MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);

2)矩阵相加：利用多个blocks

// Kernel definition__global__ void  MatAdd(floatA[N][N], floatB[N][N], floatC[N][N]) {    int  i = blockIdx.x * blockDim.x + threadIdx.x;    int  j = blockIdx.y * blockDim.y + threadIdx.y;    if(i < N && j < N)   C[i][j] = A[i][j] + B[i][j];}int  main() {    ...   // Kernel invocation  dim3  threadsPerBlock(16, 16);  dim3  numBlocks( N / threadsPerBlock.x, N / threadsPerBlock.y);  MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);

3)矩阵相乘：

计算Ｍ×Ｎ＝Ｐ：

利用一个thread block中的多个thread来计算矩阵p，每个thread负责计算矩阵p中的一个元素，步骤为：

(1)载入矩阵Ｍ的一行

(2)载入矩阵Ｎ的一列

(3)对于矩阵Ｍ的每一行的元素和矩阵Ｎ的每一列的对应元素做先相乘再相加运算。

注意：由于每一个thread block中的thread的数量不能超过1024个，所以矩阵的大小最大不能超过1024。

// Matrix multiplication kernel – thread specification__global__ void MatrixMulKernel(Matrix M, Matrix N, Matrix P) {    // 2D Thread ID    int tx = threadIdx.x;    int ty = threadIdx.y;    float z = 0; // accumulator for P     for (int k = 0; k < W; ++k)    {          z += M [ ty * W + k ]  *  N[ k * W + tx ];    }        // Write z to device memory;    P [ ty * W + tx ] = z;}

Memory model

Threads权限

-R/W per-thread registers

-R/W per-thread local memory

-R/W per-block shared memory

-R/W per-grid global memory

-Read only per-grid constant memory

-Read only per-grid texture memory

Host权限：

Ｒ/W global, constant and texture memory

对于per-block shared memory:

每个Thread block中的threads都可以共享部分block中的local memory. 对于local memory的访问速度大大高于global memory。

Shared memory的最大容量为48k。

Caffe的CUDA部分简介

Caffe中的GPU支持基于以下两点：

>1)SynchedMemory:同步内存
 CPU 和 GPU之间的内存切换是透明的。

2)每个layer的GPU实现
ConvolutionLayer::Forward_gpu( )
ConvolutionLayer::Backward_gpu( )

SynchedMemory

SyncedMemory类定义在syncedmem.hpp/cpp里, 负责caffe底层的内存管理.

内存分配与释放

内存分配与释放由两个(不属于SyncedMemory类)的内联函数完成. 代码简单直观: 如果是CPU模式, 那么调用malloc和free来申请/释放内存, 否则调用CUDA的cudaMallocHost和cudaFreeHost来申请/释放显存.

// ------ 分配内存 ------ inline void CaffeMallocHost(void** ptr, size_t size, bool* use_cuda) {#ifndef CPU_ONLY  if (Caffe::mode() == Caffe::GPU) {    CUDA_CHECK(cudaMallocHost(ptr, size));    *use_cuda = true;    return;  }#endif  *ptr = malloc(size);  *use_cuda = false;  CHECK(*ptr) << "host allocation of size " << size << " failed";}// ------ 释放内存 ------ inline void CaffeFreeHost(void* ptr, bool use_cuda) {#ifndef CPU_ONLY  if (use_cuda) {    CUDA_CHECK(cudaFreeHost(ptr));    return;  }#endif  free(ptr);}

类成员变量

  void* cpu_ptr_;  // cpu 内存地址  void* gpu_ptr_;  // gpu 内存地址  size_t size_;  // 数据大小  SyncedHead head_;  // 当前数据同步状态  bool own_cpu_data_;  //  是否是自己的cpu data? (例如set_cpu_data就是false)  bool cpu_malloc_use_cuda_;  bool own_gpu_data_;  // 是否已经申请gpu内存空间  int gpu_device_;  // 

值得稍加注意的是SyncedHead head_. 该变量的作用会在数据同步部分说明.

get and set 方法

cpu_data, gpu_data或者mutable_cpu_data, mutable_gpu_data方法返回cpu或者gpu内存指针, 前者是const void*, 不可对返回内存进行修改; 后者为void*, 可以修改.

set方法比较特别, 方法参数是指向另一段内存空间的地址:

void SyncedMemory::set_cpu_data(void* data) {  CHECK(data);  if (own_cpu_data_) {    CaffeFreeHost(cpu_ptr_, cpu_malloc_use_cuda_);  }  cpu_ptr_ = data;  head_ = HEAD_AT_CPU;  own_cpu_data_ = false;}

该函数首先释放自己申请的内存空间, 然后直接指向参数传入的内存空间 (并不是重新申请空间, 并copy数据). 最后将 own_cpu_data_设置为false, 表示外来数据(?).

保持数据同步

在调用cpu_data或者gpu_data方法时, 需要确保cpu, gpu数据内容是一致的. 这里用到了前面提到的枚举类型来记录当前同步状态

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

以to_cpu()方法为例: 检查head_所处状态, 若UNINITIALIZED, 则分配内存空间(置0); 若HEAD_AT_GPU, 则需要从GPU内存同步数据到CPU;HEAD_AT_CPU, 则说明目前最新的数据是在CPU的, 无须进行任何操作 (虽然并不知道GPU的数据是否和CPU一致, 因为当前我们并不关心GPU数据); 若SYNCED, 则CPU/GPU数据一致, 无须进行任何操作.

inline void SyncedMemory::to_cpu() {  switch (head_) {  case UNINITIALIZED:    CaffeMallocHost(&cpu_ptr_, size_, &cpu_malloc_use_cuda_);    caffe_memset(size_, 0, cpu_ptr_);    head_ = HEAD_AT_CPU;    own_cpu_data_ = true;    break;  case HEAD_AT_GPU:#ifndef CPU_ONLY    if (cpu_ptr_ == NULL) {      CaffeMallocHost(&cpu_ptr_, size_, &cpu_malloc_use_cuda_);      own_cpu_data_ = true;    }    caffe_gpu_memcpy(size_, gpu_ptr_, cpu_ptr_);    head_ = SYNCED;#else    NO_GPU;#endif    break;  case HEAD_AT_CPU:  case SYNCED:    break;  }}

Forward_gpu();

以ConvolutionalLayer中的Forward_gpu()为例：

void ConvolutionLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,  vector<Blob<Dtype>*>* top) {   const Dtype* bottom_data = bottom[0]->gpu_data();   Dtype* top_data = (*top)[0]->mutable_gpu_data();   Dtype* col_data = col_buffer_.mutable_gpu_data();   const Dtype* weight = this->blobs_[0]->gpu_data();   int weight_offset = M_ * K_;   int col_offset = K_ * N_;   int top_offset = M_ * N_;   for (int n = 0; n < NUM_; ++n) {           im2col_gpu( …);           for (int g = 0; g < GROUP_; ++g)                   caffe_gpu_gemm<Dtype>(..);   }   if (biasterm_)           caffe_gpu_gemm<Dtype>(); }