diff --git a/doc/design/ops/images/2_level_rnn.dot b/doc/design/ops/images/2_level_rnn.dot new file mode 100644 index 0000000000000000000000000000000000000000..a498e882a3d85a33d44dbad7474fa2a340e33976 --- /dev/null +++ b/doc/design/ops/images/2_level_rnn.dot @@ -0,0 +1,56 @@ +digraph G { + + rnn [label="1-th level RNN" shape=box] + + subgraph cluster0 { + label = "time step 0" + + sent0 [label="sentence"] + sent1 [label="sentence"] + + rnn1 [label="2-th level RNN" shape=box] + + sent0 -> rnn1 + sent1 -> rnn1 + } + + subgraph cluster1 { + label = "time step 1" + + sent2 [label="sentence"] + sent3 [label="sentence"] + + rnn2 [label="2-th level RNN" shape=box] + + sent2 -> rnn2 + sent3 -> rnn2 + } + + subgraph cluster2 { + label = "time step 2" + + sent4 [label="sentence"] + sent5 [label="sentence"] + + rnn3 [label="2-th level RNN" shape=box] + + sent4 -> rnn3 + sent5 -> rnn3 + } + + + para0 [label="paragraph info 0"] + para1 [label="paragraph info 1"] + para2 [label="paragraph info 2"] + + rnn1 -> para0 + rnn2 -> para1 + rnn3 -> para2 + + para0 -> rnn + para1 -> rnn + para2 -> rnn + + chapter [label="chapter info"] + rnn -> chapter +} diff --git a/doc/design/ops/images/2_level_rnn.png b/doc/design/ops/images/2_level_rnn.png new file mode 100644 index 0000000000000000000000000000000000000000..0537a75beb175c0c284717421f7aa908da2a5038 Binary files /dev/null and b/doc/design/ops/images/2_level_rnn.png differ diff --git a/doc/design/ops/images/rnn.dot b/doc/design/ops/images/rnn.dot new file mode 100644 index 0000000000000000000000000000000000000000..c1141cd9c981bb3cbf50d8bf7a6ed210280d79a5 --- /dev/null +++ b/doc/design/ops/images/rnn.dot @@ -0,0 +1,87 @@ +digraph G { + label = "simple RNN implementation" + + ranksep=2; + + //graph [nodesep=1, ranksep=1]; + + node[nodesep=1] + + subgraph cluster0 { + label = "global scope" + rankdir = TB + W + boot_memory + input + output + } + + subgraph cluster1 { + label = "step-scope 0" + rankdir = TB + memory0[label="memory"] + prememory0[label="pre-memory"] + step_input0[label="step input"] + step_output0[label="step output"] + } + + subgraph cluster2 { + label = "step-scope 1" + rankdir = TB + memory1[label="memory"] + prememory1[label="pre-memory"] + step_input1[label="step input"] + step_output1[label="step output"] + } + + subgraph cluster3 { + label = "step-scope 2" + rankdir = TB + memory2[label="memory"] + prememory2[label="pre-memory"] + step_input2[label="step input"] + step_output2[label="step output"] + } + + stepnet [shape=box] + stepnet0 [shape=box, style=dashed] + stepnet1 [shape=box, style=dashed] + stepnet2 [shape=box, style=dashed] + + + edge[color=blue] + boot_memory -> prememory0 [label="init" color="blue"] + memory0 -> prememory1 [label="copy/reference" color="blue"] + memory1 -> prememory2 [label="copy/reference" color="blue"] + + edge[color=black] + W -> stepnet0[constraint=false, style=dashed] + W -> stepnet1[constraint=false, style=dashed] + W -> stepnet2[constraint=false, style=dashed] + + memory0 -> stepnet0[style=dashed] + prememory0 -> stepnet0 -> step_output0[style=dashed] + + memory1 -> stepnet1[style=dashed] + prememory1 -> stepnet1 -> step_output1[style=dashed] + + memory2 -> stepnet2[style=dashed] + prememory2 -> stepnet2 -> step_output2[style=dashed] + + input -> step_input0 + input -> step_input1 + input -> step_input2 + + step_input0 -> stepnet0 [style=dashed] + step_input1 -> stepnet1[style=dashed] + step_input2 -> stepnet2[style=dashed] + + step_output0 -> output + step_output1 -> output + step_output2 -> output + + stepnet0 -> stepnet[style=dashed] + stepnet1 -> stepnet[style=dashed] + stepnet2 -> stepnet[style=dashed] + +} diff --git a/doc/design/ops/images/rnn.jpg b/doc/design/ops/images/rnn.jpg new file mode 100644 index 0000000000000000000000000000000000000000..9867e404cf959df0dce6ded5222b466c788fb840 Binary files /dev/null and b/doc/design/ops/images/rnn.jpg differ diff --git a/doc/design/ops/images/rnn.png b/doc/design/ops/images/rnn.png new file mode 100644 index 0000000000000000000000000000000000000000..e139e373fe8396782044cfd936fdde624f8c66fe Binary files /dev/null and b/doc/design/ops/images/rnn.png differ diff --git a/doc/design/ops/images/rnn_2level_data.dot b/doc/design/ops/images/rnn_2level_data.dot new file mode 100644 index 0000000000000000000000000000000000000000..1d85ae2617a915ad0ad8288d848b607cc37ad297 --- /dev/null +++ b/doc/design/ops/images/rnn_2level_data.dot @@ -0,0 +1,75 @@ +digraph G { + chapter [label="chapter"] + + subgraph cluster0 { + label = "paragraph 0" + + top_rnn0[label="top rnn step 0" shape=box] + + p0 [label="paragraph 0"] + p1 [label="paragraph 1"] + } + + subgraph cluster1{ + label = "paragraph 1" + + top_rnn1[label="top rnn step 1" shape=box] + + p2 [label="paragraph 0"] + p3 [label="paragraph 1"] + } + + subgraph cluster_p0 { + label = "sentence 0" + + low_rnn0 [label="low rnn step 0" shape=box] + s00 [label="sentence 0"] + s01 [label="sentence 1"] + + low_rnn0 -> s00 + low_rnn0 -> s01 + } + + subgraph cluster_p1 { + label = "sentence 1" + low_rnn1 [label="low rnn step 1" shape=box] + s10 [label="sentence 0"] + s11 [label="sentence 1"] + low_rnn1 -> s10 + low_rnn1 -> s11 + } + + subgraph cluster_p2 { + label = "sentence 1" + low_rnn2 [label="low rnn step 0" shape=box] + s20 [label="sentence 0"] + s21 [label="sentence 1"] + low_rnn2 -> s20 + low_rnn2 -> s21 + } + + subgraph cluster_p3 { + label = "sentence 1" + low_rnn3 [label="low rnn step 1" shape=box] + s30 [label="sentence 0"] + s31 [label="sentence 1"] + low_rnn3 -> s30 + low_rnn3 -> s31 + } + + + chapter -> top_rnn0 + chapter -> top_rnn1 + + top_rnn0 -> p0 + top_rnn0 -> p1 + top_rnn1 -> p2 + top_rnn1 -> p3 + + + p0 -> low_rnn0 + p1 -> low_rnn1 + p2 -> low_rnn2 + p3 -> low_rnn3 + +} diff --git a/doc/design/ops/images/rnn_2level_data.png b/doc/design/ops/images/rnn_2level_data.png new file mode 100644 index 0000000000000000000000000000000000000000..4be81b2430717a6a506342a09fc26899568574c6 Binary files /dev/null and b/doc/design/ops/images/rnn_2level_data.png differ diff --git a/doc/design/ops/rnn.md b/doc/design/ops/rnn.md new file mode 100644 index 0000000000000000000000000000000000000000..a78eea7d45e9e9553d153170aa31da55ec6e8289 --- /dev/null +++ b/doc/design/ops/rnn.md @@ -0,0 +1,153 @@ +# RNNOp design + +This document is about an RNN operator which requires that instances in a mini-batch have the same length. We will have a more flexible RNN operator. + +## RNN Algorithm Implementation + +

+ +

+ +The above diagram shows an RNN unrolled into a full network. + +There are several important concepts: + +- *step-net*: the sub-graph to run at each step, +- *memory*, $h_t$, the state of the current step, +- *ex-memory*, $h_{t-1}$, the state of the previous step, +- *initial memory value*, the ex-memory of the first step. + +### Step-scope + +There could be local variables defined in step-nets. PaddlePaddle runtime realizes these variables in *step-scopes* -- scopes created for each step. + +

+
+Figure 2 the RNN's data flow +

+ +Please be aware that all steps run the same step-net. Each step + +1. creates the step-scope, +2. realizes local variables, including step-outputs, in the step-scope, and +3. runs the step-net, which could use these variables. + +The RNN operator will compose its output from step outputs in step scopes. + +### Memory and Ex-memory + +Let's give more details about memory and ex-memory via a simply example: + +$$ +h_t = U h_{t-1} + W x_t +$$, + +where $h_t$ and $h_{t-1}$ are the memory and ex-memory of step $t$'s respectively. + +In the implementation, we can make an ex-memory variable either "refers to" the memory variable of the previous step, +or copy the value of the previous memory value to the current ex-memory variable. + +### Usage in Python + +For more information on Block, please refer to the [design doc](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/block.md). + +We can define an RNN's step-net using Block: + +```python +import paddle as pd + +X = some_op() # x is some operator's output, and is a LoDTensor +a = some_op() + +# declare parameters +W = pd.Variable(shape=[20, 30]) +U = pd.Variable(shape=[20, 30]) + +rnn = pd.create_rnn_op(output_num=1) +with rnn.stepnet(): + x = rnn.add_input(X) + # declare a memory (rnn's step) + h = rnn.add_memory(init=a) + # h.pre_state() means previous memory of rnn + new_state = pd.add_two( pd.matmul(W, x) + pd.matmul(U, h.pre_state())) + # update current memory + h.update(new_state) + # indicate that h variables in all step scopes should be merged + rnn.add_outputs(h) + +out = rnn() +``` + +Python API functions in above example: + +- `rnn.add_input` indicates the parameter is a variable that will be segmented into step-inputs. +- `rnn.add_memory` creates a variable used as the memory. +- `rnn.add_outputs` mark the variables that will be concatenated across steps into the RNN output. + +### Nested RNN and LoDTensor + +An RNN whose step-net includes other RNN operators is known as an *nested RNN*. + +For example, we could have a 2-level RNN, where the top level corresponds to paragraphs, and the lower level corresponds to sentences. + +The following figure illustrates the feeding of text into the lower level, one sentence each step, and the feeding of step outputs to the top level. The final top level output is about the whole text. + +

+ +

+ +```python +import paddle as pd + +W = pd.Variable(shape=[20, 30]) +U = pd.Variable(shape=[20, 30]) + +W0 = pd.Variable(shape=[20, 30]) +U0 = pd.Variable(shape=[20, 30]) + +# a is output of some op +a = some_op() + +# chapter_data is a set of 128-dim word vectors +# the first level of LoD is sentence +# the second level of LoD is chapter +chapter_data = pd.Variable(shape=[None, 128], type=pd.lod_tensor, level=2) + +def lower_level_rnn(paragraph): + ''' + x: the input + ''' + rnn = pd.create_rnn_op(output_num=1) + with rnn.stepnet(): + sentence = rnn.add_input(paragraph, level=0) + h = rnn.add_memory(shape=[20, 30]) + h.update( + pd.matmul(W, sentence) + pd.matmul(U, h.pre_state())) + # get the last state as sentence's info + rnn.add_outputs(h) + return rnn + +top_level_rnn = pd.create_rnn_op(output_num=1) +with top_level_rnn.stepnet(): + paragraph_data = rnn.add_input(chapter_data, level=1) + low_rnn = lower_level_rnn(paragraph_data) + paragraph_out = low_rnn() + + h = rnn.add_memory(init=a) + h.update( + pd.matmul(W0, paragraph_data) + pd.matmul(U0, h.pre_state())) + top_level_rnn.add_outputs(h) + +# just output the last step +chapter_out = top_level_rnn(output_all_steps=False) +``` + +in above example, the construction of the `top_level_rnn` calls `lower_level_rnn`. The input is a LoD Tensor. The top level RNN segments input text data into paragraphs, and the lower level RNN segments each paragraph into sentences. + +By default, the `RNNOp` will concatenate the outputs from all the time steps, +if the `output_all_steps` set to False, it will only output the final time step. + + +

+ +

diff --git a/paddle/operators/accuracy_op.cc b/paddle/operators/accuracy_op.cc new file mode 100644 index 0000000000000000000000000000000000000000..9ca04d402879b6a955d849a32175194df82b65c8 --- /dev/null +++ b/paddle/operators/accuracy_op.cc @@ -0,0 +1,66 @@ +/* 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 "paddle/operators/accuracy_op.h" + +namespace paddle { +namespace operators { + +class AccuracyOp : public framework::OperatorWithKernel { + public: + using framework::OperatorWithKernel::OperatorWithKernel; + + protected: + void InferShape(const framework::InferShapeContext &ctx) const override { + PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Inference"), + "Input of Inference must be initialized."); + PADDLE_ENFORCE_NOT_NULL(ctx.InputVar("Label"), + "Input of Inference must be initialized."); + auto *inference = ctx.Input("Inference"); + auto *label = ctx.Input("Label"); + + PADDLE_ENFORCE_EQ(label->dims().size(), 1, "label must be a vector"); + PADDLE_ENFORCE_EQ(inference->dims()[0], label->dims()[0], + "inference size must be the same as label size"); + + ctx.Output("Accuracy")->Resize({1}); + } +}; + +class AccuracyOpMaker : public framework::OpProtoAndCheckerMaker { + public: + AccuracyOpMaker(framework::OpProto *proto, + framework::OpAttrChecker *op_checker) + : OpProtoAndCheckerMaker(proto, op_checker) { + // TODO(typhoonzero): support both inference value and indices. + AddInput("Inference", "topk(indices) the network output"); + AddInput("Label", "Label of the training data"); + // TODO(typhoonzero): AddInput("Weight", ... + AddOutput("Accuracy", "The accuracy of current batch"); + + AddComment( + R"DOC(Accuracy. It will print accuracy rate for classification. +The accuracy is: +.. math:: +accuracy = \\frac{NumOfCorrectPredicts}{NumOfAllSamples})DOC"); + } +}; + +} // namespace operators +} // namespace paddle + +namespace ops = paddle::operators; +REGISTER_OP_WITHOUT_GRADIENT(accuracy, ops::AccuracyOp, ops::AccuracyOpMaker); +REGISTER_OP_CPU_KERNEL(accuracy, + ops::AccuracyKernel); diff --git a/paddle/operators/accuracy_op.cu b/paddle/operators/accuracy_op.cu new file mode 100644 index 0000000000000000000000000000000000000000..4e6d1ef9654012ce6355cbd7561c4fdc1785c11a --- /dev/null +++ b/paddle/operators/accuracy_op.cu @@ -0,0 +1,69 @@ +/* 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 "paddle/operators/accuracy_op.h" + +namespace paddle { +namespace operators { + +__global__ void AccuracySingleKernel(const int N, const int D, const int top_k, + const int* Xdata, const int* labelData, + float* accuracy) { + int correct = 0; + for (int row = 0; row < N; row++) { + const int label = labelData[row]; + for (int col = 0; col < D; col++) { + const int pred = Xdata[row * D + col]; + if (pred == label) { + ++correct; + break; + } + } + } + *accuracy = static_cast(correct) / static_cast(N); +} + +template +class AccuracyOpCUDAKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& ctx) const override { + PADDLE_ENFORCE(platform::is_gpu_place(ctx.GetPlace()), + "It must use GPUPlace."); + auto* inference = ctx.Input("Inference"); + auto* label = ctx.Input("Label"); + auto* accuracy = ctx.Output("Accuracy"); + // FIXME(typhoonzero): only support indices currently + // if add support for output values, how to detect the data type? + const int* inference_data = inference->data(); + const int* label_data = label->data(); + float* accuracy_data = accuracy->mutable_data(ctx.GetPlace()); + + size_t num_samples = inference->dims()[0]; + size_t infer_width = inference->dims()[1]; + cudaMemset((void**)&accuracy_data, 0, sizeof(float)); + + if (num_samples == 0) { + return; + } + + AccuracySingleKernel<<<1, 1>>>(num_samples, infer_width, 1, inference_data, + label_data, accuracy_data); + } +}; + +} // namespace operators +} // namespace paddle + +REGISTER_OP_GPU_KERNEL(accuracy, + paddle::operators::AccuracyOpCUDAKernel); diff --git a/paddle/operators/accuracy_op.h b/paddle/operators/accuracy_op.h new file mode 100644 index 0000000000000000000000000000000000000000..fe704efe1c979f4fc6a5a37184e51b416f5e517f --- /dev/null +++ b/paddle/operators/accuracy_op.h @@ -0,0 +1,77 @@ +/* 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. */ + +#pragma once +#include +#include "paddle/framework/eigen.h" +#include "paddle/framework/op_registry.h" + +namespace paddle { +namespace operators { + +using Tensor = framework::Tensor; + +template +using EigenMatrix = framework::EigenMatrix; + +template +using EigenVector = framework::EigenVector; + +template +using EigenScalar = framework::EigenScalar; + +template +class AccuracyKernel : public framework::OpKernel { + public: + void Compute(const framework::ExecutionContext& ctx) const override { + auto* inference = ctx.Input("Inference"); + auto* label = ctx.Input("Label"); + auto* accuracy = ctx.Output("Accuracy"); + + float* accuracy_data = accuracy->mutable_data(ctx.GetPlace()); + + const T* inference_data = inference->data(); + const T* label_data = label->data(); + + size_t num_samples = inference->dims()[0]; + size_t class_dim = inference->dims()[1]; + *accuracy_data = 0.0f; + + if (num_samples == 0) { + return; + } + + int num_correct = 0; + // assume inference is already the topk of the output + for (size_t i = 0; i < num_samples; ++i) { + PADDLE_ENFORCE_GE(label_data[i], 0, "label must >= 0"); + for (size_t j = 0; j < class_dim; ++j) { + if (inference_data[i * class_dim + j] == label_data[i]) { + ++num_correct; + break; + } + } + } + + // FIXME(typhoonzero): we don't accumulate the accuracy for now. + *accuracy_data = + static_cast(num_correct) / static_cast(num_samples); + } +}; + +} // namespace operators +} // namespace paddle diff --git a/paddle/operators/crop_op.cc b/paddle/operators/crop_op.cc index 1c048d3a70022449329fb38b3cf70efef288dc62..0c2fd096546a72507df6b6a4fea7ddfb93369b89 100644 --- a/paddle/operators/crop_op.cc +++ b/paddle/operators/crop_op.cc @@ -98,7 +98,8 @@ Out = [[1, 2], AddAttr>("shape", "A list describing the shape of output." "The size of shape list should be as same as " - "dimension size of input X."); + "dimension size of input X.") + .SetDefault(std::vector()); } }; @@ -113,8 +114,9 @@ class CropOpGrad : public framework::OperatorWithKernel { "Input(Out@GRAD) should not be null"); auto x_dims = ctx.Input("X")->dims(); auto *x_grad = ctx.Output(framework::GradVarName("X")); - - x_grad->Resize(x_dims); + if (x_grad != nullptr) { + x_grad->Resize(x_dims); + } } }; diff --git a/paddle/operators/crop_op.cu b/paddle/operators/crop_op.cu index 7977a3fe60162aed58cb2796485b554196aa39b6..a1184ce7a67d3f3089931afebe43f170f3327d25 100644 --- a/paddle/operators/crop_op.cu +++ b/paddle/operators/crop_op.cu @@ -13,12 +13,9 @@ limitations under the License. */ #define EIGEN_USE_GPU +#include #include "paddle/operators/crop_op.h" -#define CUDA_1D_KERNEL_LOOP(i, n) \ - for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; \ - i += blockDim.x * gridDim.x) - namespace paddle { namespace operators { @@ -28,39 +25,52 @@ template __global__ void CropKernel(const int N, const int64_t* out_shape, const int64_t* x_shape, const int* crop_rules, const T* x_data, T* out_data) { - CUDA_1D_KERNEL_LOOP(index, N) { - // int64_t dim_size = out_shape.size(); - int64_t pos[D]; - + int64_t pos[D]; + int tmp; + int64_t x_index; + for (int out_index = blockIdx.x * blockDim.x + threadIdx.x; out_index < N; + out_index += blockDim.x * gridDim.x) { + tmp = out_index; for (int64_t i = D - 1; i >= 0; --i) { - pos[i] = (index % out_shape[i]) + crop_rules[i * 2]; - index = index / out_shape[i]; + pos[i] = (tmp % out_shape[i]) + crop_rules[i * 2]; + tmp = tmp / out_shape[i]; } - int64_t result = pos[0]; + x_index = pos[0]; for (size_t i = 1; i < D; ++i) { - result = result * x_shape[i] + pos[i]; + x_index = x_index * x_shape[i] + pos[i]; } - - out_data[index] = x_data[result]; + out_data[out_index] = x_data[x_index]; } } template void CropCUDAFunctoin(const framework::ExecutionContext& context) { + PADDLE_ENFORCE(platform::is_gpu_place(context.GetPlace()), + "It must use GPUPlace."); auto* x = context.Input("X"); auto* out = context.Output("Out"); auto x_data = x->data(); - T* out_data = out->mutable_data(paddle::platform::CPUPlace()); + T* out_data = out->mutable_data(paddle::platform::GPUPlace()); auto x_dims = x->dims(); auto out_dims = out->dims(); int64_t out_count = framework::product(out_dims); - int64_t* x_shape = &(framework::vectorize(x_dims))[0]; - int64_t* out_shape = &(framework::vectorize(out_dims))[0]; - + int64_t x_shape[D]; + int64_t out_shape[D]; + for (int i = 0; i < D; ++i) { + x_shape[i] = x_dims[i]; + out_shape[i] = out_dims[i]; + } + int64_t* x_shape_gpu; + int64_t* out_shape_gpu; + cudaMalloc((void**)&x_shape_gpu, sizeof(int64_t) * D); + cudaMemcpy(x_shape_gpu, x_shape, sizeof(int64_t) * D, cudaMemcpyHostToDevice); + cudaMalloc((void**)&out_shape_gpu, sizeof(int64_t) * D); + cudaMemcpy(out_shape_gpu, out_shape, sizeof(int64_t) * D, + cudaMemcpyHostToDevice); auto offsets = context.op().Attr>("offsets"); PADDLE_ENFORCE_EQ( - x_dims.size(), offsets.size(), + D, offsets.size(), "Offsets size should be equal to dimension size of input tensor."); int crop_rules[D * 2]; @@ -69,13 +79,20 @@ void CropCUDAFunctoin(const framework::ExecutionContext& context) { crop_rules[i * 2 + 1] = x_dims[i] - out_dims[i] - offsets[i]; } + int* crop_rules_gpu; + cudaMalloc((void**)&crop_rules_gpu, sizeof(int) * D * 2); + cudaMemcpy(crop_rules_gpu, crop_rules, sizeof(int) * D * 2, + cudaMemcpyHostToDevice); + int n = out_dims[0]; int d = out_dims[1]; int block = 512; int grid = (n * d + block - 1) / block; - - CropKernel<<>>(out_count, out_shape, x_shape, crop_rules, - x_data, out_data); + CropKernel<<>>(out_count, out_shape_gpu, x_shape_gpu, + crop_rules_gpu, x_data, out_data); + cudaFree(crop_rules_gpu); + cudaFree(x_shape_gpu); + cudaFree(out_shape_gpu); } template diff --git a/paddle/operators/crop_op.h b/paddle/operators/crop_op.h index ff1d7694dc10e63e8c484baf58a00223264ce2b0..40bd02467492f778cd65955e077220321fca13fa 100644 --- a/paddle/operators/crop_op.h +++ b/paddle/operators/crop_op.h @@ -30,22 +30,21 @@ template void CropGradFunction(const framework::ExecutionContext& context) { auto* d_out = context.Input(framework::GradVarName("Out")); auto* d_x = context.Output(framework::GradVarName("X")); - d_x->mutable_data(context.GetPlace()); - auto d_x_dims = d_x->dims(); - auto d_out_dims = d_out->dims(); - - auto offsets = context.op().Attr>("offsets"); - - Eigen::array, D> paddings; - for (int i = 0; i < d_out_dims.size(); ++i) { - paddings[i].first = offsets[i]; - paddings[i].second = d_x_dims[i] - d_out_dims[i] - offsets[i]; + if (d_x != nullptr) { + d_x->mutable_data(context.GetPlace()); + auto d_x_dims = d_x->dims(); + auto d_out_dims = d_out->dims(); + auto offsets = context.op().Attr>("offsets"); + Eigen::array, D> paddings; + for (int i = 0; i < d_out_dims.size(); ++i) { + paddings[i].first = offsets[i]; + paddings[i].second = d_x_dims[i] - d_out_dims[i] - offsets[i]; + } + auto d_x_tensor = EigenTensor::From(*d_x); + auto d_out_tensor = EigenTensor::From(*d_out); + auto place = context.GetEigenDevice(); + d_x_tensor.device(place) = d_out_tensor.pad(paddings, 0); } - - auto d_x_tensor = EigenTensor::From(*d_x); - auto d_out_tensor = EigenTensor::From(*d_out); - auto place = context.GetEigenDevice(); - d_x_tensor.device(place) = d_out_tensor.pad(paddings, 0); } template diff --git a/paddle/operators/name_convention.md b/paddle/operators/name_convention.md index a090e0b5450509affdd739f63df618595f204f97..379385dc5d914101c7b5c9494f9383b6cf6a9b79 100644 --- a/paddle/operators/name_convention.md +++ b/paddle/operators/name_convention.md @@ -38,9 +38,11 @@ public: AccumulateOpMaker(framework::OpProto *proto, framework::OpAttrChecker *op_checker) : OpProtoAndCheckerMaker(proto, op_checker) { - AddInput("X", "(Tensor) The input tensor that has to be accumulated to the output tensor. If the output size is not the same as input size, the output tensor is first reshaped and initialized to zero, and only then, accumulation is done."); + AddInput("X", "(Tensor) The input tensor that has to be accumulated to the output tensor. + If the output size is not the same as input size, + the output tensor is first reshaped and initialized to zero, and only then, accumulation is done."); AddOutput("Out", "(Tensor) Accumulated output tensor"); - AddAttr("gamma", "(float, default 1.0) Accumulation multiplier"); + AddAttr("gamma", "(float, default 1.0) Accumulation multiplier").SetDefault(1.0f); AddComment(R"DOC( Accumulate operator accumulates the input tensor to the output tensor. If the output tensor already has the right size, we add to it; otherwise, we first @@ -51,7 +53,7 @@ Accumulation is done as shown: Out = 1*X + gamma*Out -where X is the input tensor, Y is the output tensor and gamma is the multiplier +where X is the input tensor, Out is the output tensor and gamma is the multiplier argument. )DOC"); } diff --git a/python/paddle/v2/framework/tests/test_accuracy_op.py b/python/paddle/v2/framework/tests/test_accuracy_op.py new file mode 100644 index 0000000000000000000000000000000000000000..43d60eb90d5edbd6944a11f7555f0291720dd2be --- /dev/null +++ b/python/paddle/v2/framework/tests/test_accuracy_op.py @@ -0,0 +1,25 @@ +import unittest +import numpy as np +from op_test import OpTest + + +class TestAccuracyOp(OpTest): + def setUp(self): + self.op_type = "accuracy" + infer = np.random.randint(0, 2, (32, 1)).astype("int") + label = np.random.randint(0, 2, (32, )).astype("int") + self.inputs = {'Inference': infer, "Label": label} + num_correct = 0 + for rowid in xrange(32): + for ele in infer[rowid]: + if ele == label[rowid]: + num_correct += 1 + break + self.outputs = {'Accuracy': [num_correct / 32.0]} + + def test_check_output(self): + self.check_output() + + +if __name__ == '__main__': + unittest.main() diff --git a/python/paddle/v2/framework/tests/test_crop_op.py b/python/paddle/v2/framework/tests/test_crop_op.py index 8aed80e4725b750833893bb32b2035b8d7184727..50e15086ac93d276926987487296cb6d305f99ce 100644 --- a/python/paddle/v2/framework/tests/test_crop_op.py +++ b/python/paddle/v2/framework/tests/test_crop_op.py @@ -1,8 +1,6 @@ import unittest import numpy as np -from paddle.v2.framework.op import Operator -from gradient_checker import GradientChecker -from op_test_util import OpTestMeta +from op_test import OpTest def crop(data, offsets, crop_shape): @@ -26,57 +24,68 @@ def crop(data, offsets, crop_shape): return np.array(result).reshape(crop_shape) -class TCropOp(OpTest): +class TestCropOp(OpTest): def setUp(self): - self.initTestCase() - self.type = "crop" - self.inputs = {'X': np.random.random(self.shape).astype("float32"), } + self.op_type = "crop" + self.crop_by_input = False self.attrs = {} + self.initTestCase() self.attrs['offsets'] = self.offsets - self.attrs['shape'] = self.crop_shape + if self.crop_by_input: + self.inputs = { + 'X': np.random.random(self.x_shape).astype("float32"), + 'Y': np.random.random(self.crop_shape).astype("float32") + } + else: + self.attrs['shape'] = self.crop_shape + self.inputs = { + 'X': np.random.random(self.x_shape).astype("float32"), + } self.outputs = { 'Out': crop(self.inputs['X'], self.offsets, self.crop_shape) } - print "input=%s" % self.inputs['X'] def initTestCase(self): - self.shape = (8, 8, 8) - self.crop_shape = [2, 2, 2] - self.offsets = [0, 0, 0] - - -#class TCase1(TCropOp): -# def initTestCase(self): -# self.shape = (16, 16, 16) -# self.crop_shape = [2, 2, 3] -# self.offsets = [1, 5, 3] - -#class TCropGradOp(GradientChecker): - -# def initTestCase(self): -# self.shape = (4, 4) -# self.crop_shape = [2, 2] -# self.offsets = [0, 0] - -# def setUp(self): -# self.initTestCase() -# self.op = Operator( -# type="crop", X="X", Out="Out", offsets=self.offsets, shape=self.crop_shape) -# self.inputs = {'X': np.random.random(self.shape).astype("float32"), } -# -# def test_normal(self): -# self.check_grad( -# self.op, self.inputs, set(["X"]), "Out", max_relative_error=0.5) - -#def test_cpu_gpu_compare(self): -# self.compare_grad(self.op, self.inputs) - -#class TestGradCase1(TestCropGradOp): - -# def initTestCase(self): -# self.shape = (16, 16) -# self.crop_shape = [8, 8] -# self.offsets = [1, 1] + self.x_shape = (8, 8) + self.crop_shape = [2, 2] + self.offsets = [1, 2] + + def test_check_output(self): + self.check_output() + + def test_check_grad_normal(self): + self.check_grad(['X'], 'Out', max_relative_error=0.006) + + +class TestCase1(TestCropOp): + def initTestCase(self): + self.x_shape = (16, 16, 16) + self.crop_shape = [2, 2, 3] + self.offsets = [1, 5, 3] + + +class TestCase2(TestCropOp): + def initTestCase(self): + self.x_shape = (4, 4) + self.crop_shape = [4, 4] + self.offsets = [0, 0] + + +class TestCase3(TestCropOp): + def initTestCase(self): + self.x_shape = (16, 16, 16) + self.crop_shape = [2, 2, 3] + self.offsets = [1, 5, 3] + self.crop_by_input = True + + +class TestCase4(TestCropOp): + def initTestCase(self): + self.x_shape = (4, 4) + self.crop_shape = [4, 4] + self.offsets = [0, 0] + self.crop_by_input = True + if __name__ == '__main__': unittest.main()