提交 a54565ea 编写于 作者: X xzl

delete mask pool interface from poolprojection

上级 8106f414
......@@ -35,8 +35,7 @@ limitations under the License. */
* @param[in] paddingW padding width.
* @param[out] tgtData output data.
* @param[in] tgtStride stride between output data samples.
* @param[out] maskData the location indices of select max data
* @param[in] withMask set true if output maskData
* @param[out] maskData the location indices of select max data.
*/
extern void hl_maxpool_forward(const int frameCnt,
const real* inputData,
......@@ -53,45 +52,7 @@ extern void hl_maxpool_forward(const int frameCnt,
const int paddingW,
real* tgtData,
const int tgtStride,
real* maskData,
bool withMask);
/**
* @brief Maximum pool forward.
*
* @param[in] frameCnt batch size of input image.
* @param[in] inputData input data.
* @param[in] channels number of channel.
* @param[in] height image height.
* @param[in] width image width.
* @param[in] pooledH output image height.
* @param[in] pooledW output image width.
* @param[in] sizeX width of pooling window.
* @param[in] sizeY height of pooling window.
* @param[in] strideH pooling stride height.
* @param[in] strideW pooling stride width.
* @param[in] paddingH padding height.
* @param[in] paddingW padding width.
* @param[out] tgtData output data.
* @param[in] tgtStride stride between output data samples.
* @param[out] maskData the location indices of select max data
* @param[in] withMask set true if output maskData
*/
extern void hl_maxpool_forward(const int frameCnt,
const real* inputData,
const int channels,
const int height,
const int width,
const int pooledH,
const int pooledW,
const int sizeX,
const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride);
real* maskData = NULL);
/**
* @brief Maximum pool backward.
......
......@@ -17,22 +17,6 @@ limitations under the License. */
#include "hl_cnn.h"
inline void hl_maxpool_forward(const int frameCnt,
const real* inputData,
const int channels,
const int height,
const int width,
const int pooledH,
const int pooledW,
const int sizeX,
const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride) {}
inline void hl_maxpool_forward(const int frameCnt,
const real* inputData,
const int channels,
......@@ -48,8 +32,7 @@ inline void hl_maxpool_forward(const int frameCnt,
const int paddingW,
real* tgtData,
const int tgtStride,
real* MaskData,
bool withMask) {}
real* MaskData = NULL) {}
inline void hl_maxpool_backward(const int frameCnt,
const real* inputData,
......
......@@ -32,8 +32,7 @@ __global__ void KeMaxPoolForward(const int nthreads,
const int offsetW,
real* tgtData,
const int tgtStride,
real* maskData,
bool withMask) {
real* maskData) {
int index = blockIdx.x * blockDim.x + threadIdx.x;
if (index < nthreads) {
int pw = index % pooledW;
......@@ -60,52 +59,12 @@ __global__ void KeMaxPoolForward(const int nthreads,
int tgtIndex =
index % (pooledW * pooledH * channels) + frameNum * tgtStride;
tgtData[tgtIndex] = maxval;
if (withMask) {
if (maskData != NULL) {
maskData[tgtIndex] = max_index;
}
}
}
void hl_maxpool_forward(const int frameCnt,
const real* inputData,
const int channels,
const int height,
const int width,
const int pooledH,
const int pooledW,
const int sizeX,
const int sizeY,
const int strideH,
const int strideW,
const int paddingH,
const int paddingW,
real* tgtData,
const int tgtStride) {
int num_kernels = pooledH * pooledW * channels * frameCnt;
int blocks = (num_kernels + 1024 - 1) / 1024;
dim3 threads(1024, 1);
dim3 grid(blocks, 1);
KeMaxPoolForward<<<grid, threads, 0, STREAM_DEFAULT>>>(num_kernels,
inputData,
channels,
height,
width,
pooledH,
pooledW,
sizeX,
sizeY,
strideH,
strideW,
paddingH,
paddingW,
tgtData,
tgtStride,
NULL,
false);
CHECK_SYNC("hl_maxpool_forward failed");
}
void hl_maxpool_forward(const int frameCnt,
const real* inputData,
const int channels,
......@@ -121,8 +80,7 @@ void hl_maxpool_forward(const int frameCnt,
const int paddingW,
real* tgtData,
const int tgtStride,
real* maskData,
bool withMask) {
real* maskData) {
int num_kernels = pooledH * pooledW * channels * frameCnt;
int blocks = (num_kernels + 1024 - 1) / 1024;
dim3 threads(1024, 1);
......@@ -143,8 +101,7 @@ void hl_maxpool_forward(const int frameCnt,
paddingW,
tgtData,
tgtStride,
maskData,
withMask);
maskData);
CHECK_SYNC("hl_maxpool_forward failed");
}
......
......@@ -13,6 +13,7 @@ See the License for the specific language governing permissions and
limitations under the License. */
#include "PoolLayer.h"
#include "MaxPoolWithMaskLayer.h"
#include "PoolProjectionLayer.h"
#include "paddle/utils/Logging.h"
#ifdef PADDLE_WITH_CUDA
......@@ -44,24 +45,20 @@ bool PoolLayer::init(const LayerMap& layerMap,
strideY_ = conf.has_stride_y() ? conf.stride_y() : conf.stride();
confPaddingY_ = conf.has_padding_y() ? conf.padding_y() : conf.padding();
outputY_ = conf.has_output_y() ? conf.output_y() : conf.output_x();
with_mask_ = false;
if (poolType_ == "max-pool-with-mask") {
setOutput("mask", &mask_);
with_mask_ = true;
}
return true;
}
Layer* PoolLayer::create(const LayerConfig& config) {
CHECK_EQ(config.inputs_size(), 1);
const std::string& pool = config.inputs(0).pool_conf().pool_type();
if (pool == "max-projection" || pool == "avg-projection" ||
pool == "max-pool-with-mask") {
if (pool == "max-projection" || pool == "avg-projection") {
return new PoolProjectionLayer(config);
#ifdef PADDLE_WITH_CUDA
} else if (CudnnPoolLayer::typeCheck(pool)) {
return new CudnnPoolLayer(config);
#endif
} else if (pool == "max-pool-with-mask") {
return new MaxPoolWithMaskLayer(config);
} else {
LOG(FATAL) << "Unknown pool type: " << pool;
return nullptr;
......
......@@ -37,8 +37,6 @@ protected:
int confPaddingY_;
std::string poolType_;
bool with_mask_;
Argument mask_;
public:
explicit PoolLayer(const LayerConfig& config) : Layer(config) {}
......
......@@ -36,10 +36,6 @@ PoolProjection::PoolProjection(const ProjectionConfig& config,
strideY_ = conf.has_stride_y() ? conf.stride_y() : conf.stride();
confPaddingY_ = conf.has_padding_y() ? conf.padding_y() : conf.padding();
outputY_ = conf.has_output_y() ? conf.output_y() : conf.output_x();
with_mask_ = false;
if (poolType_ == "max-pool-with-mask") {
with_mask_ = true;
}
}
size_t PoolProjection::getSize() {
......@@ -77,8 +73,6 @@ PoolProjection* PoolProjection::create(const ProjectionConfig& config,
return new MaxPoolProjection(config, parameter, useGpu);
} else if (pool == "avg-projection") {
return new AvgPoolProjection(config, parameter, useGpu);
} else if (pool == "max-pool-with-mask") {
return new MaxPoolProjection(config, parameter, useGpu);
} else {
LOG(FATAL) << "Unknown pool type: " << pool;
return nullptr;
......@@ -90,10 +84,7 @@ void MaxPoolProjection::forward() {
CHECK_EQ(width, out_->value->getWidth());
MatrixPtr inputV = in_->value;
MatrixPtr outV = out_->value;
MatrixPtr maskV = out_->value;
if (with_mask_) {
maskV = mask_->value;
}
outV->maxPoolForward(*inputV,
imgSizeY_,
imgSize_,
......@@ -105,9 +96,7 @@ void MaxPoolProjection::forward() {
outputY_,
outputX_,
confPaddingY_,
confPadding_,
maskV,
with_mask_);
confPadding_);
}
void MaxPoolProjection::backward(const UpdateCallback& callback) {
......@@ -180,26 +169,4 @@ void AvgPoolProjection::backward(const UpdateCallback& callback) {
confPaddingY_,
confPadding_);
}
void MaxWithMaskPoolProjection::forward() {
size_t width = getSize();
CHECK_EQ(width, out_->value->getWidth());
MatrixPtr inputV = in_->value;
MatrixPtr outV = out_->value;
MatrixPtr maskV = mask_->value;
outV->maxPoolForward(*inputV,
imgSizeY_,
imgSize_,
channels_,
sizeX_,
sizeY_,
strideY_,
stride_,
outputY_,
outputX_,
confPaddingY_,
confPadding_,
maskV,
with_mask_);
}
} // namespace paddle
......@@ -28,7 +28,6 @@ protected:
int confPaddingY_, confPadding_;
size_t channels_;
std::string poolType_;
bool with_mask_;
public:
PoolProjection(const ProjectionConfig& config,
......@@ -65,14 +64,4 @@ public:
virtual void backward(const UpdateCallback& callback = nullptr);
};
class MaxWithMaskPoolProjection : public MaxPoolProjection {
public:
MaxWithMaskPoolProjection(const ProjectionConfig& config,
ParameterPtr parameter,
bool useGpu)
: MaxPoolProjection(config, parameter, useGpu) {}
virtual void forward();
};
} // namespace paddle
......@@ -52,15 +52,8 @@ void PoolProjectionLayer::forward(PassType passType) {
int batchSize = in.value->getHeight();
int size = getSize();
if (with_mask_) {
resetSpecifyOutput(mask_,
batchSize,
size,
/* isValueClean */ false,
/* isGradClean */ true);
}
resetOutput(batchSize, size);
poolProjection_->forward(&in, &output_, &mask_, passType);
poolProjection_->forward(&in, &output_, passType);
}
void PoolProjectionLayer::backward(const UpdateCallback& callback) {
......
......@@ -69,17 +69,6 @@ public:
forward();
}
void forward(const Argument* in,
const Argument* out,
const Argument* mask,
PassType passType) {
in_ = in;
out_ = out;
mask_ = mask;
passType_ = passType;
forward();
}
virtual void prefetch(const Argument* in) {}
virtual void forward() = 0;
virtual void backward(const UpdateCallback& callback) = 0;
......@@ -141,8 +130,6 @@ protected:
const Argument* in_;
/// Store `out` passed to forward()
const Argument* out_;
/// Store `mask` passed to forward()
const Argument* mask_;
/// Store `passType` passed to forward()
PassType passType_;
/// Layer forward function
......
......@@ -68,7 +68,7 @@ void doOneMaxPoolingWithMaskOutputTest(MatrixPtr& inputMat,
std::vector<DataLayerPtr> dataLayers;
LayerMap layerMap;
vector<Argument> datas;
;
initDataLayer(config,
&dataLayers,
&datas,
......@@ -85,7 +85,7 @@ void doOneMaxPoolingWithMaskOutputTest(MatrixPtr& inputMat,
LayerPtr maxPoolingWithMaskOutputLayer;
initTestLayer(config, &layerMap, &parameters, &maxPoolingWithMaskOutputLayer);
maxPoolingWithMaskOutputLayer->forward(PASS_GC);
;
checkMatrixEqual(maxPoolingWithMaskOutputLayer->getOutput("mask").value,
maskMat);
}
......@@ -105,13 +105,15 @@ TEST(Layer, maxPoolingWithMaskOutputLayerFwd) {
maskMat->setData(maskData);
doOneMaxPoolingWithMaskOutputTest(
inputMat, "max-pool-with-mask", useGpu, maskMat);
#ifdef PADDLE_WITH_CUDA
useGpu = true;
inputMat = Matrix::create(1, 25, false, useGpu);
maskMat = Matrix::create(1, 4, false, useGpu);
inputMat->copyFrom(inputData, 25);
maskMat->copyFrom(maskData, 4);
doOneMaxPoolingWithMaskOutputTest(
inputMat, "max-pool-with-mask", useGpu, maskMat);
#endif
/*
#ifdef PADDLE_WITH_CUDA
useGpu = true;
inputMat = Matrix::create(1, 25, false, useGpu);
maskMat = Matrix::create(1, 4, false, useGpu);
inputMat->copyFrom(inputData, 25);
maskMat->copyFrom(maskData, 4);
doOneMaxPoolingWithMaskOutputTest(
inputMat, "max-pool-with-mask", useGpu, maskMat);
#endif
*/
}
......@@ -1017,34 +1017,6 @@ void GpuMatrix::check(std::ostream& os, Matrix& refMat, bool printDiff) {
LOG(INFO) << "the diffCnt is " << diffCnt;
}
void GpuMatrix::maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
size_t channels,
size_t sizeX,
size_t sizeY,
size_t strideH,
size_t strideW,
size_t outputH,
size_t outputW,
size_t paddingH,
size_t paddingW) {
maxPoolForward(inputMat,
imgSizeH,
imgSizeW,
channels,
sizeX,
sizeY,
strideH,
strideW,
outputH,
outputW,
paddingH,
paddingW,
NULL,
false);
}
void GpuMatrix::maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
......@@ -1057,8 +1029,7 @@ void GpuMatrix::maxPoolForward(Matrix& inputMat,
size_t outputW,
size_t paddingH,
size_t paddingW,
MatrixPtr maskMatP,
bool withMask) {
MatrixPtr maskMatP) {
CHECK(inputMat.useGpu_ == true) << "Matrix type are not equal";
real* inputData = inputMat.getData();
......@@ -1068,7 +1039,7 @@ void GpuMatrix::maxPoolForward(Matrix& inputMat,
CHECK(height_ == inputMat.getHeight());
CHECK(width_ == outputH * outputW * channels);
if (withMask) {
if (maskMatP != NULL) {
CHECK(maskMatP->useGpu_ == true) << "Matrix type are not equal";
CHECK(outputH * outputW * channels == maskMatP->getWidth());
maskData = maskMatP->getData();
......@@ -1089,8 +1060,7 @@ void GpuMatrix::maxPoolForward(Matrix& inputMat,
paddingW,
data_,
getStride(),
maskData,
withMask);
maskData);
}
void GpuMatrix::maxPoolBackward(Matrix& inputMat,
......@@ -2001,34 +1971,6 @@ void CpuMatrix::inverse(MatrixPtr& matInv, bool memAlloc) {
CHECK_EQ(info, 0);
}
void CpuMatrix::maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
size_t channels,
size_t sizeX,
size_t sizeY,
size_t strideH,
size_t strideW,
size_t outputH,
size_t outputW,
size_t paddingH,
size_t paddingW) {
maxPoolForward(inputMat,
imgSizeH,
imgSizeW,
channels,
sizeX,
sizeY,
strideH,
strideW,
outputH,
outputW,
paddingH,
paddingW,
NULL,
false);
}
void CpuMatrix::maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
......@@ -2041,8 +1983,7 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
size_t outputW,
size_t paddingH,
size_t paddingW,
MatrixPtr maskMatP,
bool withMask) {
MatrixPtr maskMatP) {
real* inputData = inputMat.getData();
real* outData = data_;
real* maskData = NULL;
......@@ -2054,7 +1995,7 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
CHECK_EQ(channels * outLength, this->getWidth());
size_t outStride = getStride();
if (withMask) {
if (maskMatP != NULL) {
maskData = maskMatP->getData();
CHECK_EQ(channels * outLength, maskMatP->getWidth());
}
......@@ -2080,7 +2021,7 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
int wstart = pw * strideW - paddingW;
int wend = std::min(wstart + sizeX, imgSizeW);
wstart = std::max(wstart, 0);
if (!withMask) {
if (maskMatP == NULL) {
for (int h = hstart; h < hend; ++h) {
for (int w = wstart; w < wend; ++w) {
outData[ph * outputW + pw] = std::max(
......@@ -2103,7 +2044,7 @@ void CpuMatrix::maxPoolForward(Matrix& inputMat,
inputData += inLength;
outData += outLength;
if (withMask) maskData += outLength;
if (maskMatP != NULL) maskData += outLength;
}
}
}
......
......@@ -861,26 +861,7 @@ public:
/**
* Pooling forward operation, pick out the largest element
* in the sizeX of value.
*/
virtual void maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
size_t channels,
size_t sizeX,
size_t sizeY,
size_t strideH,
size_t strideW,
size_t outputH,
size_t outputW,
size_t paddingH,
size_t paddingW) {
LOG(FATAL) << "Not implemeted";
}
/**
* Pooling forward operation, pick out the largest element
* in the sizeX of value, if set withMask true, it will
* in the sizeX of value, if the maskMatP is not NULL, it will
* also caculate the location indices.
*/
virtual void maxPoolForward(Matrix& inputMat,
......@@ -895,8 +876,7 @@ public:
size_t outputW,
size_t paddingH,
size_t paddingW,
MatrixPtr maskMatP,
bool withMask) {
MatrixPtr maskMatP = NULL) {
LOG(FATAL) << "Not implemeted";
}
......@@ -1437,19 +1417,6 @@ public:
void classificationError(Matrix& output, IVector& label, size_t topkSize = 1);
void maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
size_t channels,
size_t sizeX,
size_t sizeY,
size_t strideH,
size_t strideW,
size_t outputH,
size_t outputW,
size_t paddingH,
size_t paddingW);
void maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
......@@ -1462,8 +1429,7 @@ public:
size_t outputW,
size_t paddingH,
size_t paddingW,
MatrixPtr maskMatP,
bool withMask);
MatrixPtr maskMatP);
void maxPoolBackward(Matrix& image,
size_t imgSizeH,
......@@ -1723,19 +1689,6 @@ public:
MatrixPtr clone(size_t height, size_t width, bool useGpu = false);
void maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
size_t channels,
size_t sizeX,
size_t sizeY,
size_t strideH,
size_t strideW,
size_t outputH,
size_t outputW,
size_t paddingH,
size_t paddingW);
void maxPoolForward(Matrix& inputMat,
size_t imgSizeH,
size_t imgSizeW,
......@@ -1748,8 +1701,7 @@ public:
size_t outputW,
size_t paddingH,
size_t paddingW,
MatrixPtr maskMatP,
bool withMask);
MatrixPtr maskMatP);
void maxPoolBackward(Matrix& image,
size_t imgSizeH,
......
Markdown is supported
0% .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑!
想要评论请 注册