关于maxpooling和avgpooling

面试题经常会出maxpooling和avgpooling的优化，具体思路是

1）对于avgpooling，使用动态规划

2）对于maxpooling，使用队列维护历史最大值

具体代码细节后续再补充。

学习了这两种方法后，有点好奇真正的框架中是怎么实现pooling的，于是去找caffe的底层代码，发现并不是这样操作。

具体操作思路：对于GPU上计算pooling，由于有大量的计算单元，可以起大量的计算线程，使用队列反而会更低效。caffe的做法是对于每一个n,c维度上的pooling单元（blob)，都使用单独的一个线程去负责实现。

CUDA_KERNEL_LOOP(index, nthreads)是开启nthreads个线程，每个index线程负责一个blob。

由于数据是一维数组存储的，对于blob(n, c, ph, pw) ，其中n,c,ph,pw为blob开始处的下标。

pw, ph, c, n的计算就是通过index计算到原feauture map(bottom_data)中的下标。计算方法为四维坐标映射到一维数组：

index = ((n * C + c) * H + ph)*W + pw = nCHW + cHW + phW + pw

其中，大写的NCHW为一个output的feature map(top data)的各个维度的最大宽度。

所以 ph = index / W - nCH - cH - pw / W，由于pw < W，ph < H

则ph = ph % H = index / W % H

其他pw, n, c同理

然后后面就是在blob内计算最大值，mask保存最大值的位置，然后backward的时候使用

放代码：

__global__ void MaxPoolForward(const int nthreads,
    const Dtype* const bottom_d