m53297601
/
Paddle
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

update interface of context projection functions, Tensor -> Matrix/Vector

Browse Source
avx_docs
xutianbing 9 years ago
parent 2c37ad7edc
commit ea4d08dab6
6 changed files with 207 additions and 181 deletions
Show Stats Download Patch File Download Diff File

185
paddle/function/ContextProjectionOp.cpp
View File

File diff suppressed because it is too large Load Diff

33
paddle/function/ContextProjectionOp.h
Unescape View File

@ -32,14 +32,13 @@ namespace paddle {
* *
*/ */
template <DeviceType Device> template <DeviceType Device>
void ContextProjectionForward(Tensor& output, void ContextProjectionForward(typename MatrixT<Device>::type* output,
const Tensor& input, const typename MatrixT<Device>::type* input,
const Tensor& weight, const typename MatrixT<Device>::type* weight,
const Tensor& sequence, const typename SequenceT<Device>::type& sequence,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t begin_pad, size_t begin_pad);
bool is_padding);
/** /**
* \brief Context Projection Backward. * \brief Context Projection Backward.
@ -55,10 +54,10 @@ void ContextProjectionForward(Tensor& output,
* *
*/ */
template <DeviceType Device> template <DeviceType Device>
void ContextProjectionBackward(Tensor& out_grad, void ContextProjectionBackward(typename MatrixT<Device>::type* out_grad,
Tensor& in_grad, typename MatrixT<Device>::type* in_grad,
Tensor& w_grad, typename MatrixT<Device>::type* w_grad,
const Tensor& sequence, const typename SequenceT<Device>::type& seq_vec,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t begin_pad, size_t begin_pad,
@ -66,16 +65,18 @@ void ContextProjectionBackward(Tensor& out_grad,
size_t total_pad); size_t total_pad);
template <DeviceType Device> template <DeviceType Device>
void ContextProjectionBackwardData(Tensor& out_grad, void ContextProjectionBackwardData(
Tensor& in_grad, typename MatrixT<Device>::type* out_grad,
const Tensor& sequence, typename MatrixT<Device>::type* in_grad,
const typename SequenceT<Device>::type& sequence,
size_t context_length, size_t context_length,
int context_start); int context_start);
template <DeviceType Device> template <DeviceType Device>
void ContextProjectionBackwardWeight(Tensor& out_grad, void ContextProjectionBackwardWeight(
Tensor& w_grad, typename MatrixT<Device>::type* out_grad,
const Tensor& sequence, typename MatrixT<Device>::type* w_grad,
const typename SequenceT<Device>::type& seq_vec,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t total_pad, size_t total_pad,

108
paddle/function/ContextProjectionOpGpu.cu
Unescape View File

@ -75,18 +75,16 @@ __global__ void KeContextProjectionForward(const real* input,
void hl_context_projection_forward(const real* input, void hl_context_projection_forward(const real* input,
const int* sequence, const int* sequence,
real* weight, const real* weight,
real* output, real* output,
int num_sequences, int num_sequences,
int input_dim, int input_dim,
int context_length, int context_length,
int context_start, int context_start,
int begin_pad, int begin_pad) {
bool is_padding) {
CHECK_NOTNULL(input); CHECK_NOTNULL(input);
CHECK_NOTNULL(sequence); CHECK_NOTNULL(sequence);
CHECK_NOTNULL(output); CHECK_NOTNULL(output);
CHECK(!is_padding || weight);
int block_size = 128; int block_size = 128;
int blocks_x = num_sequences; int blocks_x = num_sequences;
@ -94,7 +92,7 @@ void hl_context_projection_forward(const real* input,
dim3 threads(block_size, 1); dim3 threads(block_size, 1);
dim3 grid(blocks_x, blocks_y); dim3 grid(blocks_x, blocks_y);
if (is_padding) { if (weight) {
KeContextProjectionForward<true><<< grid, threads, 0, STREAM_DEFAULT >>> KeContextProjectionForward<true><<< grid, threads, 0, STREAM_DEFAULT >>>
(input, sequence, weight, output, input_dim, (input, sequence, weight, output, input_dim,
context_length, context_start, begin_pad); context_length, context_start, begin_pad);
@ -107,31 +105,23 @@ void hl_context_projection_forward(const real* input,
} }
template <> template <>
void ContextProjectionForward<DEVICE_TYPE_GPU>(Tensor& output, void ContextProjectionForward<DEVICE_TYPE_GPU>(GpuMatrix* output,
const Tensor& input, const GpuMatrix* input,
const Tensor& weight, const GpuMatrix* weight,
const Tensor& sequence, const GpuIVector& sequence,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t begin_pad, size_t begin_pad) {
bool is_padding) { CHECK(input && output);
CHECK(output.getData() && input.getData() && sequence.getData()); hl_context_projection_forward(input->getData(),
CHECK_EQ(output.dims_.size(), 2); sequence.getData(),
CHECK_EQ(input.dims_.size(), 2); weight ? weight->getData() : nullptr,
CHECK_EQ(weight.dims_.size(), 2); output->getData(),
CHECK_EQ(sequence.dims_.size(), 1); sequence.getSize() - 1,
CHECK_EQ(output.dims_[1], input.dims_[1] * context_length); input->getWidth(),
hl_context_projection_forward(input.getData(),
reinterpret_cast<int*>(sequence.getData()),
weight.getData(),
output.getData(),
sequence.dims_[0] - 1,
input.dims_[1],
context_length, context_length,
context_start, context_start,
begin_pad, begin_pad);
is_padding);
} }
__global__ void KeContextProjectionBackwardData(real* out_grad, __global__ void KeContextProjectionBackwardData(real* out_grad,
@ -200,22 +190,17 @@ void hl_context_projection_backward_data(real* out_grad,
} }
template <> template <>
void ContextProjectionBackwardData<DEVICE_TYPE_GPU>(Tensor& out_grad, void ContextProjectionBackwardData<DEVICE_TYPE_GPU>(GpuMatrix* out_grad,
Tensor& in_grad, GpuMatrix* in_grad,
const Tensor& sequence, const GpuIVector& sequence,
size_t context_length, size_t context_length,
int context_start) { int context_start) {
CHECK(in_grad.getData() && out_grad.getData() && sequence.getData()); CHECK(in_grad && out_grad);
CHECK_EQ(out_grad.dims_.size(), 2); hl_context_projection_backward_data(out_grad->getData(),
CHECK_EQ(in_grad.dims_.size(), 2); sequence.getData(),
CHECK_EQ(sequence.dims_.size(), 1); in_grad->getData(),
CHECK_EQ(out_grad.dims_[1], in_grad.dims_[1] * context_length); sequence.getSize() - 1,
in_grad->getWidth(),
hl_context_projection_backward_data(out_grad.getData(),
reinterpret_cast<int*>(sequence.getData()),
in_grad.getData(),
sequence.dims_[0] - 1,
in_grad.dims_[1],
context_length, context_length,
context_start); context_start);
} }
@ -320,24 +305,20 @@ void hl_context_projection_backward_weight(real* out_grad,
} }
template <> template <>
void ContextProjectionBackwardWeight<DEVICE_TYPE_GPU>(Tensor& out_grad, void ContextProjectionBackwardWeight<DEVICE_TYPE_GPU>(
Tensor& w_grad, GpuMatrix* out_grad,
const Tensor& sequence, GpuMatrix* w_grad,
const GpuIVector& seq_vec,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t total_pad, size_t total_pad,
size_t begin_pad) { size_t begin_pad) {
CHECK(w_grad.getData() && out_grad.getData() && sequence.getData()); CHECK(out_grad && w_grad);
CHECK_EQ(out_grad.dims_.size(), 2); hl_context_projection_backward_weight(out_grad->getData(),
CHECK_EQ(w_grad.dims_.size(), 2); seq_vec.getData(),
CHECK_EQ(sequence.dims_.size(), 1); w_grad->getData(),
CHECK_EQ(out_grad.dims_[1], w_grad.dims_[1] * context_length); seq_vec.getSize() - 1,
w_grad->getWidth(),
hl_context_projection_backward_weight(out_grad.getData(),
reinterpret_cast<int*>(sequence.getData()),
w_grad.getData(),
sequence.dims_[0] - 1,
w_grad.dims_[1],
total_pad, total_pad,
context_length, context_length,
context_start, context_start,
@ -345,24 +326,27 @@ void ContextProjectionBackwardWeight<DEVICE_TYPE_GPU>(Tensor& out_grad,
} }
template <> template <>
void ContextProjectionBackward<DEVICE_TYPE_GPU>(Tensor& out_grad, void ContextProjectionBackward<DEVICE_TYPE_GPU>(GpuMatrix* out_grad,
Tensor& in_grad, GpuMatrix* in_grad,
Tensor& w_grad, GpuMatrix* w_grad,
const Tensor& sequence, const GpuIVector& sequence,
size_t context_length, size_t context_length,
int context_start, int context_start,
size_t begin_pad, size_t begin_pad,
bool is_padding, bool is_padding,
size_t total_pad) { size_t total_pad) {
if (in_grad.getData()) { CHECK(out_grad);
ContextProjectionBackwardData<DEVICE_TYPE_GPU>(out_grad, if (in_grad) {
ContextProjectionBackwardData<DEVICE_TYPE_GPU>(
out_grad,
in_grad, in_grad,
sequence, sequence,
context_length, context_length,
context_start); context_start);
} }
if (is_padding && w_grad.getData()) { if (is_padding && w_grad) {
ContextProjectionBackwardWeight<DEVICE_TYPE_GPU>(out_grad, ContextProjectionBackwardWeight<DEVICE_TYPE_GPU>(
out_grad,
w_grad, w_grad,
sequence, sequence,
context_length, context_length,

3
paddle/function/ContextProjectionOpTest.cpp
Unescape View File

@ -32,8 +32,7 @@ void testMatrixProjectionForward(int context_start,
FuncConfig() FuncConfig()
.set("context_length", context_length) .set("context_length", context_length)
.set("context_start", context_start) .set("context_start", context_start)
.set("begin_pad", std::max(0, -context_start)) .set("begin_pad", std::max(0, -context_start)));
.set("is_padding", is_padding));
CpuMatrix cpu_in(batch_size, input_dim); CpuMatrix cpu_in(batch_size, input_dim);
cpu_in.randomizeUniform(); cpu_in.randomizeUniform();

13
paddle/function/Function.h
Unescape View File

@ -40,6 +40,19 @@ struct MatrixT<DEVICE_TYPE_GPU> {
using type = GpuMatrix; using type = GpuMatrix;
}; };
template <DeviceType Device>
struct SequenceT;
template <>
struct SequenceT<DEVICE_TYPE_CPU> {
using type = CpuIVector;
};
template <>
struct SequenceT<DEVICE_TYPE_GPU> {
using type = GpuIVector;
};
typedef std::vector<size_t> Dims; typedef std::vector<size_t> Dims;
class Tensor { class Tensor {

12
paddle/gserver/layers/ContextProjection.cpp
Unescape View File

@ -53,8 +53,7 @@ bool ContextProjection::init() {
FuncConfig() FuncConfig()
.set("context_length", context_length) .set("context_length", context_length)
.set("context_start", context_start) .set("context_start", context_start)
.set("begin_pad", beginPad_) .set("begin_pad", beginPad_));
.set("is_padding", is_padding));
createFunction(backward_, createFunction(backward_,
"ContextProjectionBackward", "ContextProjectionBackward",
FuncConfig() FuncConfig()
@ -112,7 +111,7 @@ void ContextProjection::forward() {
size_t dim = out_->value->getWidth(); size_t dim = out_->value->getWidth();
CHECK_EQ(dim, input_dim * config_.context_length()); CHECK_EQ(dim, input_dim * config_.context_length());
size_t batch_size = in_->value->getHeight(); size_t batch_size = in_->value->getHeight();
CHECK_EQ(batch_size, out_->value->getHeight()); CHECK_EQ(forward_.size(), 1) << "Only one forward function here";
REGISTER_TIMER_INFO("ContextProjectionForward", getName().c_str()); REGISTER_TIMER_INFO("ContextProjectionForward", getName().c_str());
bool is_padding = config_.trainable_padding(); bool is_padding = config_.trainable_padding();
@ -120,12 +119,6 @@ void ContextProjection::forward() {
auto w_ptr = auto w_ptr =
state_ ? state_.get() : is_padding ? weight_->getW().get() : nullptr; state_ ? state_.get() : is_padding ? weight_->getW().get() : nullptr;
auto start_pos = in_->sequenceStartPositions; auto start_pos = in_->sequenceStartPositions;
/// if use state_ as weight_, w_ptr already has mem, so padding true
forward_[0]->init(FuncConfig()
.set("context_length", config_.context_length())
.set("context_start", config_.context_start())
.set("begin_pad", beginPad_)
.set("is_padding", state_ ? true : is_padding));
forward_[0]->calc({Tensor(in_->value->getData(), Dims{batch_size, input_dim}), forward_[0]->calc({Tensor(in_->value->getData(), Dims{batch_size, input_dim}),
Tensor(w_ptr ? w_ptr->getData() : nullptr, Tensor(w_ptr ? w_ptr->getData() : nullptr,
Dims{w_ptr ? w_ptr->getHeight() : 0, input_dim}), Dims{w_ptr ? w_ptr->getHeight() : 0, input_dim}),
@ -161,6 +154,7 @@ void ContextProjection::backward(const UpdateCallback& callback) {
CHECK_EQ(dim, input_dim * config_.context_length()); CHECK_EQ(dim, input_dim * config_.context_length());
size_t batch_size = in_->value->getHeight(); size_t batch_size = in_->value->getHeight();
CHECK_EQ(batch_size, out_->value->getHeight()); CHECK_EQ(batch_size, out_->value->getHeight());
CHECK_EQ(backward_.size(), 1) << "Only one backward function here";
REGISTER_TIMER_INFO("ContextProjectionBackward", getName().c_str()); REGISTER_TIMER_INFO("ContextProjectionBackward", getName().c_str());
bool is_padding = config_.trainable_padding(); bool is_padding = config_.trainable_padding();

Write Preview
Loading…
Cancel
Save