maxpoolingとavgpoolingについて


面接の問題はいつもmaxpoolingとavgpoolingの最適化を出して、具体的な構想は
1)avgpoolingの場合、動的計画を使用する
2)maxpoolingの場合、キューを使用して履歴最大値を維持する
具体的なコードの詳細は後で補足します.
この2つの方法を学んだ後、本当のフレームワークの中でどのようにpoolingを実現したのか少し好奇心を持って、caffeの下位コードを探して、このような操作ではないことを発見しました.
具体的な操作の構想:GPUの上でpoolingを計算して、大量の計算のユニットがあるため、大量の計算のスレッドを果たすことができて、キューを使うのはかえって更に低い効果があります.caffeのやり方は,n,c次元のpoolingユニット(blob)ごとに,単独のスレッドを用いて実現を担当することである.
CUDA_KERNEL_LOOP(index,nthreads)は、nthreads個のスレッドを開き、各index個のスレッドがblobを担当する.
データは1次元配列で格納されるため、n,c,ph,pwはblob(n,c,ph,pw)の先頭にある下付き記号である.
pw,ph,c,nの計算はindexによって元のfeauture map(bottom_data)に計算された下付きスケールです.計算方法は4次元座標を1次元配列にマッピングします.
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はpwph=ph%H=index/W%H
その他pw,n,c同理
その後、blob内で最大値を計算し、maskが最大値を保存する場所、backwardの場合に使用します.
コード:
__global__ void MaxPoolForward(const int nthreads,
    const Dtype* const bottom_data, const int num, const int channels,
    const int height, const int width, const int pooled_height,
    const int pooled_width, const int kernel_h, const int kernel_w,
    const int stride_h, const int stride_w, const int pad_h, const int pad_w,
    Dtype* const top_data, int* mask, Dtype* top_mask) {
  CUDA_KERNEL_LOOP(index, nthreads) {
    const int pw = index % pooled_width;
    const int ph = (index / pooled_width) % pooled_height;
    const int c = (index / pooled_width / pooled_height) % channels;
    const int n = index / pooled_width / pooled_height / channels;
    int hstart = ph * stride_h - pad_h;
    int wstart = pw * stride_w - pad_w;
    const int hend = min(hstart + kernel_h, height);
    const int wend = min(wstart + kernel_w, width);
    hstart = max(hstart, 0);
    wstart = max(wstart, 0);
    Dtype maxval = -FLT_MAX;
    int maxidx = -1;
    const Dtype* const bottom_slice =
        bottom_data + (n * channels + c) * height * width;
    for (int h = hstart; h < hend; ++h) {
      for (int w = wstart; w < wend; ++w) {
        if (bottom_slice[h * width + w] > maxval) {
          maxidx = h * width + w;
          maxval = bottom_slice[maxidx];
        }
      }
    }
    top_data[index] = maxval;
    if (mask) {
      mask[index] = maxidx;
    } else {
      top_mask[index] = maxidx;
    }
  }
}