Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into lookup_table
dangqingqing
8 years ago
commit
1795e57671
@ -0,0 +1,107 @@
|
|||||||
|
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||||
|
|
||||||
|
Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
you may not use this file except in compliance with the License.
|
||||||
|
You may obtain a copy of the License at
|
||||||
|
|
||||||
|
http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
|
||||||
|
Unless required by applicable law or agreed to in writing, software
|
||||||
|
distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
See the License for the specific language governing permissions and
|
||||||
|
limitations under the License. */
|
||||||
|
|
||||||
|
#include "Layer.h"
|
||||||
|
|
||||||
|
namespace paddle {
|
||||||
|
|
||||||
|
/**
|
||||||
|
* A layer applies a linear transformation to each element in each row of
|
||||||
|
* the input matrix. For each element, the layer first re-scale it and then
|
||||||
|
* adds a bias to it.
|
||||||
|
*
|
||||||
|
* \f[
|
||||||
|
* y = wx + b
|
||||||
|
* \f]
|
||||||
|
*
|
||||||
|
* Here, w is the scale and b is the bias. Both w and b are trainable scalars.
|
||||||
|
*
|
||||||
|
*/
|
||||||
|
|
||||||
|
class ScaleShiftLayer : public Layer {
|
||||||
|
protected:
|
||||||
|
std::unique_ptr<Weight> scale_;
|
||||||
|
std::unique_ptr<Weight> offset_;
|
||||||
|
|
||||||
|
public:
|
||||||
|
explicit ScaleShiftLayer(const LayerConfig& config) : Layer(config) {}
|
||||||
|
|
||||||
|
bool init(const LayerMap& layerMap,
|
||||||
|
const ParameterMap& parameterMap) override;
|
||||||
|
|
||||||
|
void forward(PassType passType) override;
|
||||||
|
void backward(const UpdateCallback& callback = nullptr) override;
|
||||||
|
};
|
||||||
|
|
||||||
|
REGISTER_LAYER(scale_shift, ScaleShiftLayer);
|
||||||
|
|
||||||
|
bool ScaleShiftLayer::init(const LayerMap& layerMap,
|
||||||
|
const ParameterMap& parameterMap) {
|
||||||
|
Layer::init(layerMap, parameterMap);
|
||||||
|
CHECK_EQ(inputLayers_.size(), 1U);
|
||||||
|
scale_.reset(new Weight(1, 1, parameters_[0]));
|
||||||
|
if (biasParameter_.get() != NULL) {
|
||||||
|
offset_ = std::unique_ptr<Weight>(new Weight(1, 1, biasParameter_));
|
||||||
|
}
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
void ScaleShiftLayer::forward(PassType passType) {
|
||||||
|
Layer::forward(passType);
|
||||||
|
|
||||||
|
MatrixPtr inV = getInputValue(0);
|
||||||
|
resetOutput(inV->getHeight(), inV->getWidth());
|
||||||
|
MatrixPtr outV = getOutputValue();
|
||||||
|
real scaleValue = scale_->getW()->getElement(0, 0);
|
||||||
|
outV->mulScalar(*inV, scaleValue);
|
||||||
|
if (offset_) {
|
||||||
|
real offsetValue = offset_->getW()->getElement(0, 0);
|
||||||
|
outV->add(offsetValue);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
void ScaleShiftLayer::backward(const UpdateCallback& callback) {
|
||||||
|
MatrixPtr inV = getInputValue(0);
|
||||||
|
MatrixPtr inG = getInputGrad(0);
|
||||||
|
MatrixPtr outV = getOutputValue();
|
||||||
|
MatrixPtr outG = getOutputGrad();
|
||||||
|
|
||||||
|
/* Calculate the parameter gradient for the current layer */
|
||||||
|
if (scale_->getWGrad()) {
|
||||||
|
MatrixPtr rowSumMtx;
|
||||||
|
Matrix::resizeOrCreate(rowSumMtx, outG->getHeight(), 1, false, useGpu_);
|
||||||
|
// this_i = scaleDest * this_i + scaleSum * \sum_j b_{ij} * c_{ij}
|
||||||
|
rowSumMtx->sumOfProducts(
|
||||||
|
/* b= */ *inV, /* c= */ *outG, /* scaleSum= */ 1, /* scaleDest= */ 0.);
|
||||||
|
// this_i = scaleDest * this_i + scaleSum * \sum_j b_{ji}
|
||||||
|
scale_->getWGrad()->sumCols(
|
||||||
|
/* b= */ *rowSumMtx, /* scaleSum= */ 1., /* scaleDest= */ 1.);
|
||||||
|
scale_->getParameterPtr()->incUpdate(callback);
|
||||||
|
}
|
||||||
|
if (offset_ && offset_->getWGrad()) {
|
||||||
|
MatrixPtr rowSumMtx;
|
||||||
|
Matrix::resizeOrCreate(rowSumMtx, outG->getHeight(), 1, false, useGpu_);
|
||||||
|
rowSumMtx->sumRows(*outG, 1., 0.);
|
||||||
|
offset_->getWGrad()->sumCols(*rowSumMtx, 1., 1.);
|
||||||
|
offset_->getParameterPtr()->incUpdate(callback);
|
||||||
|
}
|
||||||
|
|
||||||
|
/* Calculate the input layers error */
|
||||||
|
if (inG) {
|
||||||
|
real scaleValue = scale_->getW()->getElement(0, 0);
|
||||||
|
inG->add(*outG, scaleValue);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
} // namespace paddle
|
||||||
@ -1,53 +1,65 @@
|
|||||||
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
/* Copyright (c) 2016 PaddlePaddle Authors. All Rights Reserve.
|
||||||
|
|
||||||
Licensed under the Apache License, Version 2.0 (the "License");
|
Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
you may not use this file except in compliance with the License.
|
you may not use this file except in compliance with the License.
|
||||||
You may obtain a copy of the License at
|
You may obtain a copy of the License at
|
||||||
|
|
||||||
http://www.apache.org/licenses/LICENSE-2.0
|
http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
|
||||||
Unless required by applicable law or agreed to in writing, software
|
Unless required by applicable law or agreed to in writing, software
|
||||||
distributed under the License is distributed on an "AS IS" BASIS,
|
distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
See the License for the specific language governing permissions and
|
See the License for the specific language governing permissions and
|
||||||
limitations under the License. */
|
limitations under the License. */
|
||||||
|
|
||||||
#include <memory>
|
#include <thrust/device_ptr.h>
|
||||||
#include <random>
|
#include <thrust/iterator/counting_iterator.h>
|
||||||
#include "paddle/platform/dynload/curand.h"
|
#include <thrust/random.h>
|
||||||
#include "paddle/platform/gpu_info.h"
|
#include <thrust/transform.h>
|
||||||
|
|
||||||
#include "paddle/framework/op_registry.h"
|
#include "paddle/framework/op_registry.h"
|
||||||
|
#include "paddle/framework/operator.h"
|
||||||
|
|
||||||
namespace paddle {
|
namespace paddle {
|
||||||
namespace operators {
|
namespace operators {
|
||||||
|
|
||||||
template <typename T>
|
template <typename T>
|
||||||
class GaussianRandomKernel : public framework::OpKernel {
|
struct GaussianGenerator {
|
||||||
|
T mean_, std_;
|
||||||
|
unsigned int seed_;
|
||||||
|
|
||||||
|
__host__ __device__ GaussianGenerator(T mean, T std, int seed)
|
||||||
|
: mean_(mean), std_(std), seed_(seed) {}
|
||||||
|
|
||||||
|
__host__ __device__ T operator()(const unsigned int n) const {
|
||||||
|
thrust::minstd_rand rng;
|
||||||
|
rng.seed(seed_);
|
||||||
|
thrust::normal_distribution<T> dist(mean_, std_);
|
||||||
|
rng.discard(n);
|
||||||
|
return dist(rng);
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
class GPUGaussianRandomKernel : public framework::OpKernel {
|
||||||
public:
|
public:
|
||||||
void Compute(const framework::ExecutionContext& context) const override {
|
void Compute(const framework::ExecutionContext& context) const override {
|
||||||
float mean = context.op_.GetAttr<float>("mean");
|
auto* tensor = context.Output<framework::Tensor>("Out");
|
||||||
float std = context.op_.GetAttr<float>("std");
|
|
||||||
auto* tensor = context.Output<framework::Tensor>(0);
|
|
||||||
T* data = tensor->mutable_data<T>(context.GetPlace());
|
T* data = tensor->mutable_data<T>(context.GetPlace());
|
||||||
|
unsigned int seed =
|
||||||
int seed = context.op_.GetAttr<int>("seed");
|
static_cast<unsigned int>(context.op_.GetAttr<int>("seed"));
|
||||||
if (seed == 0) {
|
if (seed == 0) {
|
||||||
std::random_device rd;
|
std::random_device rd;
|
||||||
seed = rd();
|
seed = rd();
|
||||||
}
|
}
|
||||||
curandGenerator_t g;
|
T mean = static_cast<T>(context.op_.GetAttr<float>("mean"));
|
||||||
PADDLE_ENFORCE(platform::dynload::curandCreateGenerator(
|
T std = static_cast<T>(context.op_.GetAttr<float>("std"));
|
||||||
&g, CURAND_RNG_PSEUDO_DEFAULT));
|
thrust::counting_iterator<unsigned int> index_sequence_begin(0);
|
||||||
PADDLE_ENFORCE(
|
ssize_t N = framework::product(tensor->dims());
|
||||||
platform::dynload::curandSetPseudoRandomGeneratorSeed(g, seed));
|
thrust::transform(index_sequence_begin, index_sequence_begin + N,
|
||||||
platform::dynload::curandGenerateNormal(
|
thrust::device_ptr<T>(data),
|
||||||
g, data, framework::product(tensor->dims()), mean, std);
|
GaussianGenerator<T>(mean, std, seed));
|
||||||
}
|
}
|
||||||
};
|
};
|
||||||
|
|
||||||
} // namespace operators
|
} // namespace operators
|
||||||
} // namespace paddle
|
} // namespace paddle
|
||||||
|
|
||||||
namespace ops = paddle::operators;
|
REGISTER_OP_GPU_KERNEL(gaussian_random,
|
||||||
REGISTER_OP_GPU_KERNEL(gaussian_random, ops::GaussianRandomKernel<float>);
|
paddle::operators::GPUGaussianRandomKernel<float>);
|
||||||
|
|||||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue