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の場合に使用します.
コード:
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はpw
その他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;
}
}
}