Skip to content

Commit 8067a42

Browse files
committed
refine avg-pooling, which is exclusive. refine related code.
1 parent 59c48f9 commit 8067a42

File tree

2 files changed

+142
-172
lines changed

2 files changed

+142
-172
lines changed

paddle/cuda/src/hl_cuda_cnn.cu

Lines changed: 13 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -211,13 +211,11 @@ __global__ void KeAvgPoolForward(const int nthreads,
211211

212212
int hstart = ph * strideH - padH;
213213
int wstart = pw * strideW - padW;
214-
int hend = min(hstart + sizeY, height + padH);
215-
int wend = min(wstart + sizeX, width + padW);
216-
int pool_size = (hend - hstart) * (wend - wstart);
214+
int hend = min(hstart + sizeY, height);
215+
int wend = min(wstart + sizeX, width);
217216
hstart = max(hstart, 0);
218217
wstart = max(wstart, 0);
219-
hend = min(hend, height);
220-
wend = min(wend, width);
218+
int pool_size = (hend - hstart) * (wend - wstart);
221219

222220
real aveval = 0;
223221
inputData += (frameNum * channels + c) * height * width;
@@ -299,11 +297,11 @@ __global__ void KeAvgPoolBackward(const int nthreads,
299297
outGrad += (frameNum * outStride + offsetC * pooledH * pooledW);
300298

301299
for (int ph = phstart; ph < phend; ++ph) {
300+
int hstart = ph * strideH - padH;
301+
int hend = min(hstart + sizeY, height + padH);
302302
for (int pw = pwstart; pw < pwend; ++pw) {
303303
// figure out the pooling size
304-
int hstart = ph * strideH - padH;
305304
int wstart = pw * strideW - padW;
306-
int hend = min(hstart + sizeY, height + padH);
307305
int wend = min(wstart + sizeX, width + padW);
308306
int poolsize = (hend - hstart) * (wend - wstart);
309307
gradient += outGrad[ph * pooledW + pw] / poolsize;
@@ -600,16 +598,13 @@ __global__ void KeAvgPool3DForward(const int nthreads,
600598
int dstart = pd * strideD - padD;
601599
int hstart = ph * strideH - padH;
602600
int wstart = pw * strideW - padW;
603-
int dend = min(dstart + sizeZ, depth + padD);
604-
int hend = min(hstart + sizeY, height + padH);
605-
int wend = min(wstart + sizeX, width + padW);
606-
int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart);
601+
int dend = min(dstart + sizeZ, depth);
602+
int hend = min(hstart + sizeY, height);
603+
int wend = min(wstart + sizeX, width);
607604
dstart = max(dstart, 0);
608605
hstart = max(hstart, 0);
609606
wstart = max(wstart, 0);
610-
dend = min(dend, depth);
611-
hend = min(hend, height);
612-
wend = min(wend, width);
607+
int pool_size = (dend - dstart) * (hend - hstart) * (wend - wstart);
613608

614609
real aveval = 0;
615610
inputData += (frameNum * channels + c) * depth * height * width;
@@ -712,14 +707,14 @@ __global__ void KeAvgPool3DBackward(const int nthreads,
712707
outGrad += (frameNum * channels + offsetC) * pooledD * pooledH * pooledW;
713708

714709
for (int pd = pdstart; pd < pdend; ++pd) {
710+
int dstart = pd * strideD - padD;
711+
int dend = min(dstart + sizeZ, depth + padD);
715712
for (int ph = phstart; ph < phend; ++ph) {
713+
int hstart = ph * strideH - padH;
714+
int hend = min(hstart + sizeY, height + padH);
716715
for (int pw = pwstart; pw < pwend; ++pw) {
717716
// figure out the pooling size
718-
int dstart = pd * strideD - padD;
719-
int hstart = ph * strideH - padH;
720717
int wstart = pw * strideW - padW;
721-
int dend = min(dstart + sizeZ, depth + padD);
722-
int hend = min(hstart + sizeY, height + padH);
723718
int wend = min(wstart + sizeX, width + padW);
724719
int poolsize = (dend - dstart) * (hend - hstart) * (wend - wstart);
725720
gradient += outGrad[(pd * pooledH + ph) * pooledW + pw] / poolsize;

0 commit comments

Comments
 (0)