Merge branch 'develop' of https://github.com/PaddlePaddle/Paddle into develop
zchen0211
8 years ago
commit
0a0f194836
@ -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>);
|
||||||
|
|||||||
Loading…
Reference in new issue