'
// where type is 0 for default, 1 for send, and 2 for receive.
// No channel and values are needed for default cases.
@@ -150,7 +150,7 @@ into **X**. It will also create a temp variable called **case_to_execute**. Th
filled in by the select_op after it has completed processing the case statements.
If there are no available cases to execute (ie: all cases are blocked on channel operations, and
-there is no default statement), then the select_op will block the current thread. The thread will
+there is no default statement), then the select_op will block the current thread. The thread will
unblock once there is a channel operation affecting one of the case statements, at which point, the
**select_op** will set the **case_to_execute** variable to the index of the case to execute.
@@ -247,17 +247,17 @@ blocks {
```
-Cases are represented by a **conditional_block operator**, whose's condition is set as the output of
-equal(**case_to_execute**, **case_index**). Since each case index is unique in this sub-block,
+Cases are represented by a **conditional_block operator**, whose's condition is set as the output of
+equal(**case_to_execute**, **case_index**). Since each case index is unique in this sub-block,
only one case will be executed.
### select_op flow
-
+
-The select algorithm is inspired by golang's select routine. Please refer to
+The select algorithm is inspired by golang's select routine. Please refer to
http://www.tapirgames.com/blog/golang-concurrent-select-implementation for more information.
## Backward Pass
diff --git a/doc/fluid/design/dist_train/distributed_architecture.md b/doc/fluid/design/dist_train/distributed_architecture.md
index a405cb6aaf..229cb47c17 100644
--- a/doc/fluid/design/dist_train/distributed_architecture.md
+++ b/doc/fluid/design/dist_train/distributed_architecture.md
@@ -40,11 +40,11 @@ computation is only specified in Python code which sits outside of PaddlePaddle,
Similar to how a compiler uses an intermediate representation (IR) so that the programmer does not need to manually optimize their code for most of the cases, we can have an intermediate representation in PaddlePaddle as well. The compiler optimizes the IR as follows:
-
+
PaddlePaddle can support model parallelism by converting the IR so that the user no longer needs to manually perform the computation and operations in the Python component:
-
+
The IR for PaddlePaddle after refactoring is called a `Block`, it specifies the computation dependency graph and the variables used in the computation.
@@ -60,7 +60,7 @@ For a detailed explanation, refer to this document -
The revamped distributed training architecture can address the above discussed limitations. Below is the illustration of how it does so:
-
+
The major components are: *Python API*, *Distribute Transpiler* and *Remote Executor*.
@@ -152,7 +152,7 @@ for data in train_reader():
`JobDesc` object describe the distributed job resource specification to run on
Cluster environment.
-
+
`RemoteExecutor.run` sends the `ProgramDesc` and
[TrainingJob](https://github.com/PaddlePaddle/cloud/blob/unreleased-tpr/doc/autoscale/README.md#training-job-resource)
@@ -171,7 +171,7 @@ In the future, a more general placement algorithm should be implemented, which m
The local training architecture will be the same as the distributed training architecture, the difference is that everything runs locally, and there is just one PaddlePaddle runtime:
-
+
### Training Data
diff --git a/doc/fluid/design/dist_train/multi_cpu.md b/doc/fluid/design/dist_train/multi_cpu.md
index a8d8ee0422..38222d0830 100644
--- a/doc/fluid/design/dist_train/multi_cpu.md
+++ b/doc/fluid/design/dist_train/multi_cpu.md
@@ -8,11 +8,11 @@ Op graph to a multi-CPU Op graph, and run `ParallelDo` Op to run the graph.
## Transpiler
-
+
After converted:
-
+
## Implement
diff --git a/doc/fluid/design/dist_train/parameter_server.md b/doc/fluid/design/dist_train/parameter_server.md
index 6ce48dfbfc..73c85da5e8 100644
--- a/doc/fluid/design/dist_train/parameter_server.md
+++ b/doc/fluid/design/dist_train/parameter_server.md
@@ -41,11 +41,11 @@ We will need these OPs: *Send*, *Recv*, *Enqueue*, *Dequeue*.
Below is an example of converting the user defined graph to the
subgraphs for the trainer and the parameter server:
-
+
After converting:
-
+
1. The parameter variable W and its optimizer program are placed on the parameter server.
1. Operators are added to the program.
@@ -69,8 +69,7 @@ In Fluid, we introduce [SelectedRows](../selected_rows.md) to represent a list o
non-zero gradient data. So when we do parameter optimization both locally and remotely,
we only need to send those non-zero rows to the optimizer operators:
-
-
+
### Benefits
- Model parallelism becomes easier to implement: it is an extension to
diff --git a/doc/fluid/design/dynamic_rnn/rnn.md b/doc/fluid/design/dynamic_rnn/rnn.md
index 6f414e5549..7b61b050f6 100644
--- a/doc/fluid/design/dynamic_rnn/rnn.md
+++ b/doc/fluid/design/dynamic_rnn/rnn.md
@@ -5,7 +5,7 @@ This document describes the RNN (Recurrent Neural Network) operator and how it i
## RNN Algorithm Implementation
-
+
The above diagram shows an RNN unrolled into a full network.
@@ -22,7 +22,7 @@ There are several important concepts here:
There could be local variables defined in each step-net. PaddlePaddle runtime realizes these variables in *step-scopes* which are created for each step.
-
+
Figure 2 illustrates the RNN's data flow
@@ -93,7 +93,7 @@ For example, we could have a 2-level RNN, where the top level corresponds to par
The following figure illustrates feeding in text into the lower level, one sentence at a step, and the feeding in step outputs to the top level. The final top level output is about the whole text.
-
+
```python
@@ -149,5 +149,5 @@ If the `output_all_steps` is set to False, it will only output the final time st
-
+
diff --git a/doc/fluid/design/modules/batch_norm_op.md b/doc/fluid/design/modules/batch_norm_op.md
index d1392619c4..e451ffcc73 100644
--- a/doc/fluid/design/modules/batch_norm_op.md
+++ b/doc/fluid/design/modules/batch_norm_op.md
@@ -2,7 +2,7 @@
## What is batch normalization
-Batch normalization is a frequently-used method in deep network training. It adjusts the mean and variance of a layer's output, and make the data distribution easier for next layer's training.
+Batch normalization is a frequently-used method in deep network training. It adjusts the mean and variance of a layer's output, and make the data distribution easier for next layer's training.
The principle of batch normalization can be summarized into a simple function:
@@ -66,7 +66,7 @@ As most C++ operators do, `batch_norm_op` is defined by inputs, outputs, attribu
The following graph showes the training computational process of `batch_norm_op`:
-
+
cudnn provides APIs to finish the whole series of computation, we can use them in our GPU kernel.
@@ -74,13 +74,13 @@ cudnn provides APIs to finish the whole series of computation, we can use them i
`batch_norm_op` is warpped as a layer in Python:
-```python
-def batch_norm_layer(net,
+```python
+def batch_norm_layer(net,
input,
- output,
- scale,
- bias,
- use_global_est = False,
+ output,
+ scale,
+ bias,
+ use_global_est = False,
epsilon = 1e-6,
momentum = 0.99):
mean_cache = scope.new_var(name = 'estimated_mean', trainable = False)
@@ -119,15 +119,15 @@ for pass_id in range(PASS_NUM):
if pass_id % 100 == 0:
net.infer(test_image) # run inferencing model
# ...
-```
+```
`is_infer` is an attribute. Once an operator is created, its attributes can not be changed. It suggests us that we shall maintain two `batch_norm_op` in the model, one's `is_infer` is `True`(we call it `infer_batch_norm_op`) and the other one's is `False`(we call it `train_batch_norm_op`). They share all parameters and variables, but be placed in two different branches. That is to say, if a network contains a `batch_norm_op`, it will fork into two branches, one go through `train_batch_norm_op` and the other one go through `infer_batch_norm_op`:
-
+
+
+
+| Python classes |
+Protobuf messages |
+
+
+
+
+| Program |
+ProgramDesc |
+
+
+| Block |
+BlockDesc |
+
+
+| Operator |
+OpDesc |
+
+
+| Variable |
+VarDesc |
+
+
+
+
Please be aware that these Python classes need to maintain some construction-time information, which are not part of the protobuf messages.
diff --git a/doc/fluid/design/modules/regularization.md b/doc/fluid/design/modules/regularization.md
index 21280ac898..8cd5ff71d1 100644
--- a/doc/fluid/design/modules/regularization.md
+++ b/doc/fluid/design/modules/regularization.md
@@ -6,23 +6,23 @@ A central problem in machine learning is how to design an algorithm that will pe
### Parameter Norm Penalties
Most common regularization approaches in deep learning are based on limiting the capacity of the models by adding a parameter norm penalty to the objective function `J`. This is given as follows:
-
+
The parameter `alpha` is a hyperparameter that weights the relative contribution of the norm penalty term, `omega`, relative to the standard objective function `J`.
The most commonly used norm penalties are the L2 norm penalty and the L1 norm penalty. These are given as follows:
##### L2 Regularization:
-
+
##### L1 Regularization
-
+
A much more detailed mathematical background of regularization can be found [here](http://www.deeplearningbook.org/contents/regularization.html).
## Regularization Survey
-A detailed survey of regularization in various deep learning frameworks can be found [here](https://github.com/PaddlePaddle/Paddle/wiki/Regularization-Survey).
+A detailed survey of regularization in various deep learning frameworks can be found [here](https://github.com/PaddlePaddle/Paddle/wiki/Regularization-Survey).
## Proposal for Regularization in PaddlePaddle
@@ -32,41 +32,35 @@ In the new design, we propose to create new operations for regularization. For n
- L2_regularization_op
- L1_regularization_op
-These ops can be like any other ops with their own CPU/GPU implementations either using Eigen or separate CPU and GPU kernels. As the initial implementation, we can implement their kernels using Eigen following the abstraction pattern implemented for [Activation Ops](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/accuracy_op.h). This abstraction pattern can make it very easy to implement new regularization schemes other than L1 and L2 norm penalties.
+These ops can be like any other ops with their own CPU/GPU implementations either using Eigen or separate CPU and GPU kernels. As the initial implementation, we can implement their kernels using Eigen following the abstraction pattern implemented for [Activation Ops](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/operators/accuracy_op.h). This abstraction pattern can make it very easy to implement new regularization schemes other than L1 and L2 norm penalties.
-The idea of building ops for regularization is in sync with the refactored Paddle philosophy of using operators to represent any computation unit. The way these ops will be added to the computation graph, will be decided by the [layer functions](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md#layer-function) in Python API.
+The idea of building ops for regularization is in sync with the refactored Paddle philosophy of using operators to represent any computation unit. The way these ops will be added to the computation graph, will be decided by the [layer functions](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md#layer-function) in Python API.
### Computation Graph
Below is an example of a really simple feed forward neural network.
-
+
The Python API will modify this computation graph to add regularization operators. The modified computation graph will look as follows:
-
+
### Python API implementation for Regularization
-Using the low level ops, `L2_regularization_op` and `L1_regularization_op`, any user can add regularization to their computation graphs. However, this will require a lot of lines of code and we should design Python APIs that support regularization. An example of such an API can be seen in [Keras](https://keras.io/regularizers/). As per the PaddlePaddle [Python API design](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md), the layer functions are responsible for creating operators, operator parameters and variables. Since regularization is a property of parameters, it makes sense to create these in the layer functions.
+Using the low level ops, `L2_regularization_op` and `L1_regularization_op`, any user can add regularization to their computation graphs. However, this will require a lot of lines of code and we should design Python APIs that support regularization. An example of such an API can be seen in [Keras](https://keras.io/regularizers/). As per the PaddlePaddle [Python API design](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md), the layer functions are responsible for creating operators, operator parameters and variables. Since regularization is a property of parameters, it makes sense to create these in the layer functions.
#### Creation of Regularization ops
There are two possibilities for creating the regularization ops:
-1. We create these ops immediately while building the computation graph.
-2. We add these ops in a lazy manner, just before the backward, similar to the way the optimization ops are added.
+1. We create these ops immediately while building the computation graph.
+2. We add these ops in a lazy manner, just before the backward, similar to the way the optimization ops are added.
-The proposal is to add these ops in a lazy manner just before the backward pass.
+The proposal is to add these ops in a lazy manner just before the backward pass.
#### Storage of Regularization attributes
-Since we want to create the regularization ops in a lazy manner, the regularization attributes (type of regularization and weight of regularization penalty) can be stored as attributes of the [`Parameter`](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/framework/framework.py#L421) class. This is because regularization is a property of the parameters and storing regularization properties with Parameters also allows for shared parameters.
+Since we want to create the regularization ops in a lazy manner, the regularization attributes (type of regularization and weight of regularization penalty) can be stored as attributes of the [`Parameter`](https://github.com/PaddlePaddle/Paddle/blob/develop/python/paddle/v2/framework/framework.py#L421) class. This is because regularization is a property of the parameters and storing regularization properties with Parameters also allows for shared parameters.
#### High-level API
In PaddlePaddle Python API, users will primarily rely on [layer functions](https://github.com/PaddlePaddle/Paddle/blob/develop/doc/design/python_api.md#layer-function) to create neural network layers. Hence, we also need to provide regularization functionality in layer functions. The design of these APIs can be postponed for later right now. A good reference for these APIs can be found in [Keras](https://keras.io/regularizers/) and also by looking at Tensorflow in [`tf.contrib.layers`](https://www.tensorflow.org/api_guides/python/contrib.layers).
-
-
-
-
-
-
diff --git a/doc/fluid/design/motivation/fluid.md b/doc/fluid/design/motivation/fluid.md
index 110b7d78bf..5e147f8263 100644
--- a/doc/fluid/design/motivation/fluid.md
+++ b/doc/fluid/design/motivation/fluid.md
@@ -10,11 +10,37 @@ Fluid is the answer. Fluid is similar to PyTorch and TensorFlow Eager Execution
Deep learning infrastructure is one of the fastest evolving technologies. Within four years, there have already been three generations of technologies invented.
-| Existed since | model as sequence of layers | model as graph of operators | No model |
-|--|--|--|--|
-| 2013 | Caffe, Theano, Torch, PaddlePaddle | | |
-| 2015 | | TensorFlow, MxNet, Caffe2, ONNX, n-graph | |
-| 2016 | | | PyTorch, TensorFlow Eager Execution, PaddlePaddle Fluid |
+
+
+
+| Existed since |
+model as sequence of layers |
+model as graph of operators |
+No model |
+
+
+
+
+| 2013 |
+Caffe, Theano, Torch, PaddlePaddle |
+ |
+ |
+
+
+| 2015 |
+ |
+TensorFlow, MxNet, Caffe2, ONNX, n-graph |
+ |
+
+
+| 2016 |
+ |
+ |
+ PyTorch, TensorFlow Eager Execution, PaddlePaddle Fluid |
+
+
+
+
From the above table, we see that the deep learning technology is evolving towards getting rid of the concept of a model. To understand the reasons behind this direction, a comparison of the *programming paradigms* or the ways to program deep learning applications using these systems, would be helpful. The following section goes over these.
diff --git a/doc/fluid/design/motivation/refactorization.md b/doc/fluid/design/motivation/refactorization.md
index 7c39fabcc6..f199cc892f 100644
--- a/doc/fluid/design/motivation/refactorization.md
+++ b/doc/fluid/design/motivation/refactorization.md
@@ -36,11 +36,37 @@ At compile time, the Python program generates a protobuf message representation
At runtime, the C++ program realizes the graph and runs it.
-| | Representation (protobuf messages) | Realization (C++ class objects) |
-|---|---|---|
-|Data|[VarDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L107)|[Variable](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/variable.h#L24)|
-|Operation|[OpDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L35)|[Operator](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/operator.h#L64)|
-|Block|BlockDesc|Block|
+
+
+
+ |
+Representation (protobuf messages) |
+Realization (C++ class objects) |
+
+
+
+
+| Data |
+
+VarDesc |
+
+Variable |
+
+
+| Operation |
+
+OpDesc |
+
+Operator |
+
+
+| Block |
+BlockDesc |
+Block |
+
+
+
+
The word *graph* is interchangeable with *block* in this document. A graph consists of computation steps and local variables similar to a C++/Java program block, or a pair of parentheses(`{` and `}`).
diff --git a/doc/fluid/design/network/deep_speech_2.md b/doc/fluid/design/network/deep_speech_2.md
index af0c6ef36f..f32a5b7e8a 100644
--- a/doc/fluid/design/network/deep_speech_2.md
+++ b/doc/fluid/design/network/deep_speech_2.md
@@ -1,4 +1,4 @@
-# DeepSpeech2 on PaddlePaddle: Design Doc
+# DeepSpeech2 on PaddlePaddle: Design Doc
We are planning to build Deep Speech 2 (DS2) \[[1](#references)\], a powerful Automatic Speech Recognition (ASR) engine, on PaddlePaddle. For the first-stage plan, we have the following short-term goals:
@@ -68,11 +68,33 @@ We roughly break down the project into 14 tasks:
Tasks parallelizable within phases:
-Roadmap | Description | Parallelizable Tasks
------------ | :------------------------------------ | :--------------------
-Phase I | Simplified model & components | *Task 1* ~ *Task 8*
-Phase II | Standard model & benchmarking & profiling | *Task 9* ~ *Task 12*
-Phase III | Documentations | *Task13* ~ *Task14*
+
+
+
+| Roadmap |
+Description |
+ Parallelizable Tasks |
+
+
+
+
+| Phase I |
+Simplified model & components |
+Task 1 ~ Task 8 |
+
+
+| Phase II |
+ Standard model & benchmarking & profiling |
+Task 9 ~ Task 12 |
+
+
+| Phase III |
+ Documentations |
+ Task13 ~ Task14 |
+
+
+
+
Issue for each task will be created later. Contributions, discussions and comments are all highly appreciated and welcomed!
@@ -94,7 +116,7 @@ The classical DS2 network contains 15 layers (from bottom to top):
- **One** CTC-loss layer
-
+
Figure 1. Archetecture of Deep Speech 2 Network.
+
+
+| Required Components |
+ PaddlePaddle Support |
+ Need to Develop |
+
+
+
+
+| Data Layer I (Spectrogram) |
+Not supported yet. |
+TBD (Task 3) |
+
+
+| Data Layer II (Transcription) |
+ paddle.data_type.integer_value_sequence |
+ - |
+
+
+| 2D Convolution Layer |
+ paddle.layer.image_conv_layer |
+ - |
+
+
+| DataType Converter (vec2seq) |
+ paddle.layer.block_expand |
+ - |
+
+
+| Bi-/Uni-directional RNNs |
+paddle.layer.recurrent_group |
+ - |
+
+
+| Row Convolution Layer |
+Not supported yet. |
+TBD (Task 4) |
+
+
+| CTC-loss Layer |
+paddle.layer.warp_ctc |
+ - |
+
+
+| Batch Normalization Layer |
+paddle.layer.batch_norm |
+ - |
+
+
+| CTC-Beam search |
+Not supported yet. |
+ TBD (Task 6) |
+
+
+
+
### Row Convolution
@@ -141,18 +208,18 @@ TODO by Assignees
### Beam Search with CTC and LM
-
+
Figure 2. Algorithm for CTC Beam Search Decoder.
- 
+ 
According to the image above, the only phase that changes the LoD is beam search.
diff --git a/doc/fluid/design/others/gan_api.md b/doc/fluid/design/others/gan_api.md
index fb41df8615..7167470088 100644
--- a/doc/fluid/design/others/gan_api.md
+++ b/doc/fluid/design/others/gan_api.md
@@ -1,24 +1,24 @@
# Design for GAN
-GAN (General Adversarial Net [https://arxiv.org/abs/1406.2661]) is an important model for unsupervised learning and widely used in many areas.
+GAN (General Adversarial Net [https://arxiv.org/abs/1406.2661]) is an important model for unsupervised learning and widely used in many areas.
It applies several important concepts in machine learning system design, including building and running subgraphs, dependency tracing, different optimizers in one executor and so forth.
In our GAN design, we wrap it as a user-friendly easily customized python API to design different models. We take the conditional DC-GAN (Unsupervised Representation Learning with Deep Convolutional Generative Adversarial Networks [https://arxiv.org/abs/1511.06434]) as an example due to its good performance on image generation.
-
+
Figure 1. The overall running logic of GAN. The black solid arrows indicate the forward pass; the green dashed arrows indicate the backward pass of generator training; the red dashed arrows indicate the backward pass of the discriminator training. The BP pass of the green (red) arrow should only update the parameters in the green (red) boxes. The diamonds indicate the data providers. d\_loss and g\_loss marked in red and green are the two targets we would like to run.
The operators, layers and functions required/optional to build a GAN demo is summarized in https://github.com/PaddlePaddle/Paddle/issues/4563.
-
+
Figure 2. Photo borrowed from the original DC-GAN paper.
-## The Conditional-GAN might be a class.
+## The Conditional-GAN might be a class.
This design we adopt the popular open source design in https://github.com/carpedm20/DCGAN-tensorflow and https://github.com/rajathkmp/DCGAN. It contains following data structure:
- DCGAN(object): which contains everything required to build a GAN model. It provides following member functions methods as API:
@@ -29,7 +29,7 @@ This design we adopt the popular open source design in https://github.com/carped
Returns a generated image.
- discriminator(image):
-Given an image, decide if it is from a real source or a fake one.
+Given an image, decide if it is from a real source or a fake one.
Returns a 0/1 binary label.
- build_model(self):
@@ -47,7 +47,7 @@ To be more detailed, we introduce our design of DCGAN as following:
```python
class DCGAN(object):
def __init__(self, y_dim=None):
-
+
# hyper parameters
self.y_dim = y_dim # conditional gan or not
self.batch_size = 100
@@ -82,18 +82,18 @@ class DCGAN(object):
# input z: the random noise
# input y: input data label (optional)
# output G_im: generated fake images
-
+
if not self.y_dim:
z = pd.layer.concat(1, [z, y])
-
+
G_h0 = pd.layer.fc(z, self.G_w0, self.G_b0)
G_h0_bn = pd.layer.batch_norm(G_h0)
G_h0_relu = pd.layer.relu(G_h0_bn)
-
+
G_h1 = pd.layer.deconv(G_h0_relu, self.G_w1, self.G_b1)
G_h1_bn = pd.layer.batch_norm(G_h1)
G_h1_relu = pd.layer.relu(G_h1_bn)
-
+
G_h2 = pd.layer.deconv(G_h1_relu, self.G_W2, self.G_b2))
G_im = pd.layer.tanh(G_im)
return G_im
@@ -111,11 +111,11 @@ class DCGAN(object):
D_h0 = pd.layer.conv2d(image, w=self.D_w0, b=self.D_b0)
D_h0_bn = pd.layer.batchnorm(h0)
D_h0_relu = pd.layer.lrelu(h0_bn)
-
+
D_h1 = pd.layer.conv2d(D_h0_relu, w=self.D_w1, b=self.D_b1)
D_h1_bn = pd.layer.batchnorm(D_h1)
D_h1_relu = pd.layer.lrelu(D_h1_bn)
-
+
D_h2 = pd.layer.fc(D_h1_relu, w=self.D_w2, b=self.D_b2)
return D_h2
```
@@ -123,7 +123,7 @@ class DCGAN(object):
### Class member function: Build the model
- Define data readers as placeholders to hold the data;
- Build generator and discriminators;
-- Define two training losses for discriminator and generator, respectively.
+- Define two training losses for discriminator and generator, respectively.
If we have execution dependency engine to back-trace all tensors, the module building our GAN model will be like this:
```python
class DCGAN(object):
@@ -133,7 +133,7 @@ class DCGAN(object):
self.images = pd.data(pd.float32, [self.batch_size, self.im_size, self.im_size])
self.faked_images = pd.data(pd.float32, [self.batch_size, self.im_size, self.im_size])
self.z = pd.data(tf.float32, [None, self.z_size])
-
+
# step 1: generate images by generator, classify real/fake images with discriminator
if self.y_dim: # if conditional GAN, includes label
self.G = self.generator(self.z, self.y)
@@ -147,12 +147,12 @@ class DCGAN(object):
# generate fake images
self.sampled = self.sampler(self.z)
self.D_f = self.discriminator(self.images)
-
+
# step 2: define the two losses
self.d_loss_real = pd.reduce_mean(pd.cross_entropy(self.D_t, np.ones(self.batch_size))
self.d_loss_fake = pd.reduce_mean(pd.cross_entropy(self.D_f, np.zeros(self.batch_size))
self.d_loss = self.d_loss_real + self.d_loss_fake
-
+
self.g_loss = pd.reduce_mean(pd.cross_entropy(self.D_f, np.ones(self.batch_szie))
```
@@ -176,7 +176,7 @@ class DCGAN(object):
self.G = self.generator(self.z)
self.D_g = self.discriminator(self.G, self.y)
self.g_loss = pd.reduce_mean(pd.cross_entropy(self.D_g, np.ones(self.batch_szie))
-
+
with pd.default_block().d_block():
if self.y_dim: # if conditional GAN, includes label
self.D_t = self.discriminator(self.images, self.y)
@@ -217,7 +217,7 @@ if __name__ == "__main__":
# load mnist data
data_X, data_y = self.load_mnist()
-
+
# Two subgraphs required!!!
with pd.block().d_block():
d_optim = pd.train.Adam(lr = .001, beta= .1)
@@ -228,7 +228,7 @@ if __name__ == "__main__":
# executor
sess = pd.executor()
-
+
# training
for epoch in xrange(10000):
for batch_id in range(N / batch_size):
@@ -239,7 +239,7 @@ if __name__ == "__main__":
batch_z = np.random.uniform(-1., 1., [batch_size, z_dim])
if batch_id % 2 == 0:
- sess.run(d_step,
+ sess.run(d_step,
feed_dict = {dcgan.images: batch_im,
dcgan.y: batch_label,
dcgan.z: batch_z})
diff --git a/doc/fluid/dev/index_cn.rst b/doc/fluid/dev/index_cn.rst
index e70bf5dff3..f627437f35 100644
--- a/doc/fluid/dev/index_cn.rst
+++ b/doc/fluid/dev/index_cn.rst
@@ -4,9 +4,9 @@
.. toctree::
:maxdepth: 1
- new_op_en.md
- new_op_kernel_en.md
- use_eigen_en.md
+ new_op_cn.md
+ new_op_kernel.md
+ use_eigen_cn.md
name_convention.md
support_new_device.md
releasing_process.md
diff --git a/doc/fluid/dev/index_en.rst b/doc/fluid/dev/index_en.rst
index f0e9afcfcc..0b65fed67a 100644
--- a/doc/fluid/dev/index_en.rst
+++ b/doc/fluid/dev/index_en.rst
@@ -5,7 +5,7 @@ Development
:maxdepth: 1
new_op_en.md
- new_op_kernel_en.md
+ new_op_kernel.md
use_eigen_en.md
name_convention.md
support_new_device.md
diff --git a/doc/fluid/dev/new_op_cn.md b/doc/fluid/dev/new_op_cn.md
index 9299658567..0c3f88d9c3 100644
--- a/doc/fluid/dev/new_op_cn.md
+++ b/doc/fluid/dev/new_op_cn.md
@@ -26,13 +26,32 @@
依据是否包含kernel，可以将Op分为两种：包含Kernel的Op和不包含kernel的Op，前者Op的定义继承自`OperatorWithKernel`，后者继承自`OperatorBase`。本教程主要介绍带Kernel的Op如何写，简单总结Op需要包含的内容如下：
-
- 内容 | 定义位置
--------------- | :----------------------
-OpProtoMake定义 | `.cc`文件，Backward Op不需要定义OpProtoMake
-Op定义 | `.cc`文件
-Kernel实现 | CPU、CUDA共享Kernel实现在`.h`文件中，否则，CPU 实现在`.cc`文件中，CUDA 实现在`.cu`文件中。
-注册Op | Op注册实现在`.cc`文件；Kernel注册CPU实现在`.cc`文件中，CUDA实现在`.cu`文件中
+
+
+
+| 内容 |
+定义位置 |
+
+
+
+
+| OpProtoMake定义 |
+`.cc`文件，Backward Op不需要定义OpProtoMake |
+
+
+| Op定义 |
+ `.cc`文件 |
+
+
+| Kernel实现 |
+ CPU、CUDA共享Kernel实现在`.h`文件中，否则，CPU 实现在`.cc`文件中，CUDA 实现在`.cu`文件中。 |
+
+
+| 注册Op |
+ Op注册实现在`.cc`文件；Kernel注册CPU实现在`.cc`文件中，CUDA实现在`.cu`文件中 |
+
+
+
实现新的op都添加至目录[paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators)下，文件命名以`*_op.h`（如有） 、 `*_op.cc` 、`*_op.cu`（如有）结尾。**系统会根据文件名自动构建op和其对应的Python扩展。**
diff --git a/doc/fluid/dev/new_op_en.md b/doc/fluid/dev/new_op_en.md
index da8b1bdd10..a566a09131 100644
--- a/doc/fluid/dev/new_op_en.md
+++ b/doc/fluid/dev/new_op_en.md
@@ -33,6 +33,33 @@ Op definition | `.cc` files
Kernel implementation | The kernel methods shared between CPU and CUDA are defined in `.h` files. CPU-specific kernels live in `.cc` files, while CUDA-specific kernels are implemented in `.cu`files.
Registering the Op | Ops are registered in `.cc` files; For Kernel registration, `.cc` files contain the CPU implementation, while `.cu` files contain the CUDA implementation.
+
+
+
+| Information |
+ Where is it defined |
+
+
+
+
+| OpProtoMake definition |
+ `.cc`files, Backward Op does not need an OpProtoMake interface. |
+
+
+| Op definition |
+ `.cc` files |
+
+
+| Kernel implementation |
+ The kernel methods shared between CPU and CUDA are defined in `.h` files. CPU-specific kernels live in `.cc` files, while CUDA-specific kernels are implemented in `.cu`files. |
+
+
+| Registering the Op |
+ Ops are registered in `.cc` files; For Kernel registration, `.cc` files contain the CPU implementation, while `.cu` files contain the CUDA implementation. |
+
+
+
+
New Operator implementations are added to the list [paddle/operators](https://github.com/PaddlePaddle/Paddle/tree/develop/paddle/operators), with file names in the format `*_op.h` (if applicable), `*_op.cc`, `*_op.cu` (if applicable).** The system will use the naming scheme to automatically build operators and their corresponding Python extensions.**
@@ -279,7 +306,7 @@ A forward operator unit test inherits `unittest.TestCase` and defines metaclass
def test_check_output(self):
self.check_output()
-
+
def test_check_grad_normal(self):
self.check_grad(['X', 'Y'], 'Out', max_relative_error=0.5)
diff --git a/doc/fluid/dev/new_op_kernel_en.md b/doc/fluid/dev/new_op_kernel.md
similarity index 100%
rename from doc/fluid/dev/new_op_kernel_en.md
rename to doc/fluid/dev/new_op_kernel.md
diff --git a/doc/fluid/dev/releasing_process.md b/doc/fluid/dev/releasing_process.md
index b978726109..c5943ccd81 100644
--- a/doc/fluid/dev/releasing_process.md
+++ b/doc/fluid/dev/releasing_process.md
@@ -37,7 +37,7 @@ PaddlePaddle每次发新的版本，遵循以下流程:
可以在此页面的"Artifacts"下拉框中找到生成的3个二进制文件，分别对应CAPI，`cp27m`和`cp27mu`的版本。然后按照上述的方法
使用`twine`工具上传即可。
-
+
* 注：CI环境使用 https://github.com/PaddlePaddle/buildtools 这里的DockerImage作为编译环境以支持更多的Linux
发型版，如果需要手动编译，也可以使用这些镜像。这些镜像也可以从 https://hub.docker.com/r/paddlepaddle/paddle_manylinux_devel/tags/ 下载得到。
@@ -66,7 +66,7 @@ PaddlePaddle开发过程使用[git-flow](http://nvie.com/posts/a-successful-git-
* 建议，开发者fork的版本库使用`develop`分支同步主版本库的`develop`分支
* 建议，开发者fork的版本库中，再基于`develop`版本fork出自己的功能分支。
* 当功能分支开发完毕后，向PaddlePaddle的主版本库提交`Pull Reuqest`，进而进行代码评审。
- * 在评审过程中，开发者修改自己的代码，可以继续在自己的功能分支提交代码。
+ * 在评审过程中，开发者修改自己的代码，可以继续在自己的功能分支提交代码。
* BugFix分支也是在开发者自己的fork版本库维护，与功能分支不同的是，BugFix分支需要分别给主版本库的`master`、`develop`与可能有的`release/版本号`分支，同时提起`Pull Request`。
@@ -78,13 +78,116 @@ PaddlePaddle开发过程使用[git-flow](http://nvie.com/posts/a-successful-git-
PaddlePaddle每次发版本首先要保证PaddlePaddle Book中所有章节功能的正确性。功能的正确性包括验证PaddlePaddle目前的`paddle_trainer`训练和纯使用`Python`训练模型正确性。
-| | 新手入门章节 | 识别数字 | 图像分类 | 词向量 | 情感分析 | 语意角色标注 | 机器翻译 | 个性化推荐 |
-| --- | --- | --- | --- | --- | --- | --- | --- | --- |
-| API.V2 + Docker + GPU | | | | | | | | |
-| API.V2 + Docker + CPU | | | | | | | | |
-| `paddle_trainer` + Docker + GPU | | | | | | | | |
-| `paddle_trainer` + Docker + CPU | | | | | | | | |
-| API.V2 + Ubuntu + GPU | | | | | | | | |
-| API.V2 + Ubuntu + CPU | | | | | | | | |
-| `paddle_trainer` + Ubuntu + GPU | | | | | | | | |
-| `paddle_trainer` + Ubuntu + CPU | | | | | | | | |
+
+
+
+ |
+新手入门章节 |
+ 识别数字 |
+ 图像分类 |
+词向量 |
+ 情感分析 |
+语意角色标注 |
+ 机器翻译 |
+个性化推荐 |
+
+
+
+
+
+| API.V2 + Docker + GPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| API.V2 + Docker + CPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| `paddle_trainer` + Docker + GPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| `paddle_trainer` + Docker + CPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| API.V2 + Ubuntu + GPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| API.V2 + Ubuntu + CPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| `paddle_trainer` + Ubuntu + GPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
+| `paddle_trainer` + Ubuntu + CPU |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+ |
+
+
+
diff --git a/doc/fluid/getstarted/concepts/save_model/model_format.md b/doc/fluid/getstarted/concepts/save_model/model_format.md
index e29129fddf..1f12ba0497 100644
--- a/doc/fluid/getstarted/concepts/save_model/model_format.md
+++ b/doc/fluid/getstarted/concepts/save_model/model_format.md
@@ -4,30 +4,70 @@
A model is an output of the training process. One complete model consists of two parts, the **topology** and the **parameters**. In order to support industrial deployment, the model format must be self-complete and must not expose any training source code.
-As a result, In PaddlePaddle, the **topology** is represented as a [ProgramDesc](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/doc/design/program.md), which describes the model structure. The **parameters** contain all the trainable weights in the model. We must support large size parameters and efficient serialization/deserialization of parameters.
+As a result, In PaddlePaddle, the **topology** is represented as a [ProgramDesc](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/doc/design/program.md), which describes the model structure. The **parameters** contain all the trainable weights in the model. We must support large size parameters and efficient serialization/deserialization of parameters.
## Implementation
-The topology is saved as a plain text in a detailed self-contain protobuf file.
+The topology is saved as a plain text in a detailed self-contain protobuf file.
The parameters are saved as a binary file. As we all know, the protobuf message has a limit of [64M size](https://developers.google.com/protocol-buffers/docs/reference/cpp/google.protobuf.io.coded_stream#CodedInputStream.SetTotalBytesLimit.details). We have done a [benchmark experiment](https://github.com/PaddlePaddle/Paddle/pull/4610), which shows that protobuf is not fit for the task.
-As a result, we design a particular format for tensor serialization. By default, an arbitrary tensor in Paddle is a [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/lod_tensor.md), and has a description information proto of [LoDTensorDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L99). We save the DescProto as the byte string header. It contains all the necessary information, such as the `dims`, and the `LoD` information in [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/paddle/framework/lod_tensor.md). A tensor stores values in a continuous memory buffer. For speed we dump the raw memory to disk and save it as the byte string content. So, the binary format of one tensor is,
+As a result, we design a particular format for tensor serialization. By default, an arbitrary tensor in Paddle is a [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/lod_tensor.md), and has a description information proto of [LoDTensorDesc](https://github.com/PaddlePaddle/Paddle/blob/develop/paddle/framework/framework.proto#L99). We save the DescProto as the byte string header. It contains all the necessary information, such as the `dims`, and the `LoD` information in [LoDTensor](https://github.com/PaddlePaddle/Paddle/blob/1c0a4c901c9fc881d120249c703b15d1c50dae7d/paddle/framework/lod_tensor.md). A tensor stores values in a continuous memory buffer. For speed we dump the raw memory to disk and save it as the byte string content. So, the binary format of one tensor is,
The table below shows a tensor's byte view in detail. Note that all the signed values are written in the little-endian format.
-|field name | type | description |
-| --- | --- | --- |
-| version | uint32_t | Version of saved file. Always 0 now. |
-| tensor desc length | uint32_t | TensorDesc(Protobuf message) length in bytes. |
-| tensor desc | void* | TensorDesc protobuf binary message |
-| tensor data | void* | Tensor's data in binary format. The length of `tensor_data` is decided by `TensorDesc.dims()` and `TensorDesc.data_type()` |
-| lod_level | uint64_t | Level of LoD |
-| length of lod[0] | uint64_t | [Optional] length of lod[0] in bytes. |
-| data of lod[0] | uint64_t* | [Optional] lod[0].data() |
-| ... | ... | ... |
-
+
+
+
+| field name |
+type |
+description |
+
+
+
+
+| version |
+ uint32_t |
+ Version of saved file. Always 0 now. |
+
+
+| tensor desc length |
+ uint32_t |
+ TensorDesc(Protobuf message) length in bytes. |
+
+
+| tensor desc |
+ void* |
+ TensorDesc protobuf binary message |
+
+
+| tensor data |
+ void* |
+ Tensor's data in binary format. The length of `tensor_data` is decided by `TensorDesc.dims()` and `TensorDesc.data_type()` |
+
+
+| lod_level |
+ uint64_t |
+ Level of LoD |
+
+
+| length of lod[0] |
+ uint64_t |
+ [Optional] length of lod[0] in bytes. |
+
+
+| data of lod[0] |
+ uint64_t* |
+ [Optional] lod[0].data() |
+
+
+| ... |
+ ... |
+ ... |
+
+
+
## Summary
diff --git a/doc/fluid/howto/cluster/fluid_cluster_train_cn.md b/doc/fluid/howto/cluster/fluid_cluster_train_cn.md
index 1b6f767869..b99b90056b 100644
--- a/doc/fluid/howto/cluster/fluid_cluster_train_cn.md
+++ b/doc/fluid/howto/cluster/fluid_cluster_train_cn.md
@@ -65,10 +65,10 @@ exit(1)
**因此，在分布式的Fluid环境中，我们有两个角色需要创建，分别是Parameter Server和Trainer。**
-### 分布式训练
+### 分布式训练
Fliud专门提供了工具[Distributed Transpiler](https://github.com/PaddlePaddle/Paddle/blob/ba65d54d9d3b41cd3c5171b00f476d4e60133ddb/doc/fluid/design/dist_train/distributed_architecture.md#distributed-transpiler)用于将单机版的训练程序转换为分布式版本的训练程序。工具背后的理念是找出程序的优化算子和梯度参数，将他们分隔为两部分，通过send/recv 操作算子进行连接,优化算子和梯度参数可以在优化器的minimize函数的返回值中获取到。
```python
-optimize_ops, params_grads = sgd_optimizer.minimize(avg_cost)
+optimize_ops, params_grads = sgd_optimizer.minimize(avg_cost)
```
将Distributed Transpiler、优化算子和梯度函数放在一个代码中如下：
```python
@@ -99,15 +99,51 @@ for pass_id in range(100):
### 分布式训练脚本运行说明
分布式任务的运行需要将表格中说明的多个参数进行赋值:
-| 参数名 | 值类型 | 说明 | 示例 |
-|:-------------|:------|:---------------------------------------|:-------------|
-| trainer_id | int | 当前训练节点的ID，训练节点ID编号为0 - n-1， n为trainers的值 | 0/1/2/3 |
-| pservers | str | parameter server 列表 | 127.0.0.1:6710,127.0.0.1:6711 |
-| trainers | int | 训练节点的总个数，>0的数字 | 4 |
-| server_endpoint | str | 当前所起的服务节点的IP:PORT | 127.0.0.1:8789 |
-| training_role | str | 节点角色， TRAINER/PSERVER | PSERVER |
-
-**注意：** ```training_role```是用来区分当前所起服务的角色的，用于训练程序中，用户可根据需要自行定义，其他参数为fluid.DistributeTranspiler的transpile函数所需要，需要在调用函数前进行定义，样例如下：
+
+
+
+| 参数名 |
+ 值类型 |
+说明 |
+ 示例 |
+
+
+
+
+| trainer_id |
+ int |
+ 当前训练节点的ID，训练节点ID编号为0 - n-1， n为trainers的值 |
+ 0/1/2/3 |
+
+
+| pservers |
+ str |
+ parameter server 列表 |
+ 127.0.0.1:6710,127.0.0.1:6711 |
+
+
+| trainers |
+int |
+ 训练节点的总个数，>0的数字 |
+ 4 |
+
+
+| server_endpoint |
+ str |
+ 当前所起的服务节点的IP:PORT |
+ 127.0.0.1:8789 |
+
+
+| training_role |
+str |
+ 节点角色， TRAINER/PSERVER |
+ PSERVER |
+
+
+
+
+
+**注意：** ```training_role```是用来区分当前所起服务的角色的，用于训练程序中，用户可根据需要自行定义，其他参数为fluid.DistributeTranspiler的transpile函数所需要，需要在调用函数前进行定义，样例如下：
```python
t = fluid.DistributeTranspiler()
diff --git a/doc/fluid/howto/optimization/cpu_profiling_cn.md b/doc/fluid/howto/optimization/cpu_profiling_cn.md
index 17f895573a..8266dec3c6 100644
--- a/doc/fluid/howto/optimization/cpu_profiling_cn.md
+++ b/doc/fluid/howto/optimization/cpu_profiling_cn.md
@@ -42,14 +42,40 @@ cprofilev -a 0.0.0.0 -p 3214 -f profile.out main.py
每一列的含义是:
-| 列名 | 含义 |
-| --- | --- |
-| ncalls | 函数的调用次数 |
-| tottime | 函数实际使用的总时间。该时间去除掉本函数调用其他函数的时间 |
-| percall | tottime的每次调用平均时间 |
-| cumtime | 函数总时间。包含这个函数调用其他函数的时间 |
-| percall | cumtime的每次调用平均时间 |
-| filename:lineno(function) | 文件名, 行号，函数名 |
+
+
+
+| 列名 |
+含义 |
+
+
+
+
+| ncalls |
+ 函数的调用次数 |
+
+
+| tottime |
+ 函数实际使用的总时间。该时间去除掉本函数调用其他函数的时间 |
+
+
+| percall |
+ tottime的每次调用平均时间 |
+
+
+| cumtime |
+ 函数总时间。包含这个函数调用其他函数的时间 |
+
+
+| percall |
+ cumtime的每次调用平均时间 |
+
+
+| filename:lineno(function) |
+ 文件名, 行号，函数名 |
+
+
+
### 寻找性能瓶颈
diff --git a/doc/fluid/howto/optimization/cpu_profiling_en.md b/doc/fluid/howto/optimization/cpu_profiling_en.md
index abe4493c17..e95556dd60 100644
--- a/doc/fluid/howto/optimization/cpu_profiling_en.md
+++ b/doc/fluid/howto/optimization/cpu_profiling_en.md
@@ -57,14 +57,40 @@ port, we will see the output like the following:
where each line corresponds to Python function, and the meaning of
each column is as follows:
-| column | meaning |
-| --- | --- |
-| ncalls | the number of calls into a function |
-| tottime | the total execution time of the function, not including the execution time of other functions called by the function |
-| percall | tottime divided by ncalls |
-| cumtime | the total execution time of the function, including the execution time of other functions being called |
-| percall | cumtime divided by ncalls |
-| filename:lineno(function) | where the function is defined |
+
+
+
+| column |
+meaning |
+
+
+
+
+| ncalls |
+ the number of calls into a function |
+
+
+| tottime |
+ the total execution time of the function, not including the execution time of other functions called by the function |
+
+
+| percall |
+ tottime divided by ncalls |
+
+
+| cumtime |
+ the total execution time of the function, including the execution time of other functions being called |
+
+
+| percall |
+ cumtime divided by ncalls |
+
+
+| filename:lineno(function) |
+ where the function is define |
+
+
+
### Identify Performance Bottlenecks
diff --git a/doc/fluid/howto/performance/profiler.md b/doc/fluid/howto/performance/profiler.md
index b20b5efdc1..ee96e7c74c 100644
--- a/doc/fluid/howto/performance/profiler.md
+++ b/doc/fluid/howto/performance/profiler.md
@@ -23,7 +23,7 @@ But how to record the time for the mixed C++ and CUDA program? There many C++ A
The overall flow is shown as the following figure.
-
+
### Event
@@ -36,10 +36,10 @@ enum EventKind {
kPopRange};
```
- kMark: only a marker without time range.
-- kPushRange: mark the starting event for time range.
+- kPushRange: mark the starting event for time range.
- kPopRange: mark the ending event for time range.
-For the CPU code, the events only need to record the current time. For the CUDA code, the [event management functions of CUDA](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html#group__CUDART__EVENT) are used. For many pieces of code, an event lists are used to record each piece.
+For the CPU code, the events only need to record the current time. For the CUDA code, the [event management functions of CUDA](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html#group__CUDART__EVENT) are used. For many pieces of code, an event lists are used to record each piece.
```c++
class Event {
@@ -66,11 +66,11 @@ struct EventList {
};
```
-As mentioned above, there is no need to record the timeline when disabling the profiler. So there is a global state to enable or disable the profiler.
+As mentioned above, there is no need to record the timeline when disabling the profiler. So there is a global state to enable or disable the profiler.
```c++
enum ProfilerState {
- kDisabled,
+ kDisabled,
kCPU,
kCUDA
};
diff --git a/doc/fluid/images/2_level_rnn.dot b/doc/fluid/images/2_level_rnn.dot
new file mode 100644
index 0000000000..5d77865061
--- /dev/null
+++ b/doc/fluid/images/2_level_rnn.dot
@@ -0,0 +1,56 @@
+digraph G {
+
+ rnn [label="1st level RNN" shape=box]
+
+ subgraph cluster0 {
+ label = "time step 0"
+
+ sent0 [label="sentence"]
+ sent1 [label="sentence"]
+
+ rnn1 [label="2nd level RNN" shape=box]
+
+ sent0 -> rnn1
+ sent1 -> rnn1
+ }
+
+ subgraph cluster1 {
+ label = "time step 1"
+
+ sent2 [label="sentence"]
+ sent3 [label="sentence"]
+
+ rnn2 [label="2nd level RNN" shape=box]
+
+ sent2 -> rnn2
+ sent3 -> rnn2
+ }
+
+ subgraph cluster2 {
+ label = "time step 2"
+
+ sent4 [label="sentence"]
+ sent5 [label="sentence"]
+
+ rnn3 [label="2nd 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/fluid/images/2_level_rnn.png b/doc/fluid/images/2_level_rnn.png
new file mode 100644
index 0000000000..0537a75beb
Binary files /dev/null and b/doc/fluid/images/2_level_rnn.png differ
diff --git a/doc/fluid/images/LOD-and-shape-changes-during-decoding.jpg b/doc/fluid/images/LOD-and-shape-changes-during-decoding.jpg
new file mode 100644
index 0000000000..8b0d90f7b9
Binary files /dev/null and b/doc/fluid/images/LOD-and-shape-changes-during-decoding.jpg differ
diff --git a/doc/fluid/images/asgd.gif b/doc/fluid/images/asgd.gif
new file mode 100644
index 0000000000..4a0da7bf6d
Binary files /dev/null and b/doc/fluid/images/asgd.gif differ
diff --git a/doc/fluid/images/batch_norm_fork.dot b/doc/fluid/images/batch_norm_fork.dot
new file mode 100644
index 0000000000..4bc47713cb
--- /dev/null
+++ b/doc/fluid/images/batch_norm_fork.dot
@@ -0,0 +1,25 @@
+digraph ImageBatchNormForkGragh {
+ subgraph cluster_before {
+ Prev [label="...", shape=plaintext];
+ Rnn [label="rnn_op", shape=box];
+ BatchNorm [label="batch_norm_op", shape=box];
+ Fc [label="fc_op", shape=box];
+ After [label="...", shape=plaintext];
+ Prev -> Rnn -> BatchNorm -> Fc -> After;
+ label="original";
+ }
+
+ subgraph cluster_after {
+ Prev2 [label="...", shape=plaintext];
+ Rnn2 [label="rnn_op", shape=box];
+ BatchNorm2_1 [label="train_batch_norm_op", shape=box];
+ BatchNorm2_2 [label="infer_batch_norm_op", shape=box];
+ Fc2_1 [label="fc_op", shape=box];
+ Fc2_2 [label="fc_op", shape=box];
+ After2_1 [label="...", shape=plaintext];
+ After2_2 [label="...", shape=plaintext];
+ Prev2 -> Rnn2 -> BatchNorm2_1 -> Fc2_1 -> After2_1;
+ Rnn2 -> BatchNorm2_2 ->Fc2_2 ->After2_2
+ label="forked";
+ }
+}
diff --git a/doc/fluid/images/batch_norm_fork.png b/doc/fluid/images/batch_norm_fork.png
new file mode 100644
index 0000000000..aded62bce5
Binary files /dev/null and b/doc/fluid/images/batch_norm_fork.png differ
diff --git a/doc/fluid/images/batch_norm_op_kernel.png b/doc/fluid/images/batch_norm_op_kernel.png
new file mode 100644
index 0000000000..a99ce81ff3
Binary files /dev/null and b/doc/fluid/images/batch_norm_op_kernel.png differ
diff --git a/doc/fluid/images/beam_search.png b/doc/fluid/images/beam_search.png
new file mode 100644
index 0000000000..7f7e35f342
Binary files /dev/null and b/doc/fluid/images/beam_search.png differ
diff --git a/doc/fluid/images/ci_build_whl.png b/doc/fluid/images/ci_build_whl.png
new file mode 100644
index 0000000000..232762b82a
Binary files /dev/null and b/doc/fluid/images/ci_build_whl.png differ
diff --git a/doc/fluid/images/compiler.graffle b/doc/fluid/images/compiler.graffle
new file mode 100644
index 0000000000..8cc678fea3
Binary files /dev/null and b/doc/fluid/images/compiler.graffle differ
diff --git a/doc/fluid/images/compiler.png b/doc/fluid/images/compiler.png
new file mode 100644
index 0000000000..65d34f841a
Binary files /dev/null and b/doc/fluid/images/compiler.png differ
diff --git a/doc/fluid/images/control_flow_graph.png b/doc/fluid/images/control_flow_graph.png
new file mode 100644
index 0000000000..3579998e58
Binary files /dev/null and b/doc/fluid/images/control_flow_graph.png differ
diff --git a/doc/fluid/images/dataflow_equations.png b/doc/fluid/images/dataflow_equations.png
new file mode 100644
index 0000000000..c10f7f69f4
Binary files /dev/null and b/doc/fluid/images/dataflow_equations.png differ
diff --git a/doc/fluid/images/dcgan.png b/doc/fluid/images/dcgan.png
new file mode 100644
index 0000000000..15e8e290a1
Binary files /dev/null and b/doc/fluid/images/dcgan.png differ
diff --git a/doc/fluid/images/deep_learning.png b/doc/fluid/images/deep_learning.png
new file mode 100644
index 0000000000..026becc4d9
Binary files /dev/null and b/doc/fluid/images/deep_learning.png differ
diff --git a/doc/fluid/images/dist-graph.graffle b/doc/fluid/images/dist-graph.graffle
new file mode 100644
index 0000000000..941399c6ce
Binary files /dev/null and b/doc/fluid/images/dist-graph.graffle differ
diff --git a/doc/fluid/images/dist-graph.png b/doc/fluid/images/dist-graph.png
new file mode 100644
index 0000000000..3546b09f1c
Binary files /dev/null and b/doc/fluid/images/dist-graph.png differ
diff --git a/doc/fluid/images/distributed_architecture.graffle b/doc/fluid/images/distributed_architecture.graffle
new file mode 100644
index 0000000000..d1b6014134
Binary files /dev/null and b/doc/fluid/images/distributed_architecture.graffle differ
diff --git a/doc/fluid/images/distributed_architecture.png b/doc/fluid/images/distributed_architecture.png
new file mode 100644
index 0000000000..29c7b0c078
Binary files /dev/null and b/doc/fluid/images/distributed_architecture.png differ
diff --git a/doc/fluid/images/ds2_network.png b/doc/fluid/images/ds2_network.png
new file mode 100644
index 0000000000..1a5b2184d4
Binary files /dev/null and b/doc/fluid/images/ds2_network.png differ
diff --git a/doc/fluid/images/feed_forward.png b/doc/fluid/images/feed_forward.png
new file mode 100644
index 0000000000..d312371a04
Binary files /dev/null and b/doc/fluid/images/feed_forward.png differ
diff --git a/doc/fluid/images/feed_forward_regularized.png b/doc/fluid/images/feed_forward_regularized.png
new file mode 100644
index 0000000000..677e99bfd9
Binary files /dev/null and b/doc/fluid/images/feed_forward_regularized.png differ
diff --git a/doc/fluid/images/fluid-compiler.graffle b/doc/fluid/images/fluid-compiler.graffle
new file mode 100644
index 0000000000..c933df2cb8
Binary files /dev/null and b/doc/fluid/images/fluid-compiler.graffle differ
diff --git a/doc/fluid/images/fluid-compiler.png b/doc/fluid/images/fluid-compiler.png
new file mode 100644
index 0000000000..1b0ffed203
Binary files /dev/null and b/doc/fluid/images/fluid-compiler.png differ
diff --git a/doc/fluid/images/graph_construction_example.bash b/doc/fluid/images/graph_construction_example.bash
new file mode 100755
index 0000000000..35e6997abd
--- /dev/null
+++ b/doc/fluid/images/graph_construction_example.bash
@@ -0,0 +1,11 @@
+cat ./graph_construction_example.dot | \
+ sed 's/color=red/color=red, style=invis/g' | \
+ sed 's/color=green/color=green, style=invis/g' | \
+ dot -Tpng > graph_construction_example_forward_only.png
+
+cat ./graph_construction_example.dot | \
+ sed 's/color=green/color=green, style=invis/g' | \
+ dot -Tpng > graph_construction_example_forward_backward.png
+
+cat ./graph_construction_example.dot | \
+ dot -Tpng > graph_construction_example_all.png
diff --git a/doc/fluid/images/graph_construction_example.dot b/doc/fluid/images/graph_construction_example.dot
new file mode 100644
index 0000000000..e115f9844b
--- /dev/null
+++ b/doc/fluid/images/graph_construction_example.dot
@@ -0,0 +1,68 @@
+digraph ImageClassificationGraph {
+ ///////// The forward part /////////
+ FeedX [label="Feed", color=blue, shape=box];
+ FeedY [label="Feed", color=blue, shape=box];
+ InitW [label="Init", color=blue, shape=diamond];
+ Initb [label="Init", color=blue, shape=diamond];
+ FC [label="FC", color=blue, shape=box];
+ MSE [label="MSE", color=blue, shape=box];
+
+ x [label="x", color=blue, shape=oval];
+ l [label="l", color=blue, shape=oval];
+ y [label="y", color=blue, shape=oval];
+ W [label="W", color=blue, shape=doublecircle];
+ b [label="b", color=blue, shape=doublecircle];
+ cost [label="cost", color=blue, shape=oval];
+
+ FeedX -> x -> FC -> y -> MSE -> cost [color=blue];
+ FeedY -> l [color=blue];
+ InitW -> W [color=blue];
+ Initb -> b [color=blue];
+ W -> FC [color=blue];
+ b -> FC [color=blue];
+ l -> MSE [color=blue];
+
+ ////////// The backward part /////////
+ MSE_Grad [label="MSE_grad", color=red, shape=box];
+ FC_Grad [label="FC_grad", color=red, shape=box];
+
+ d_cost [label="d cost", color=red, shape=oval];
+ d_y [label="d y", color=red, shape=oval];
+ d_b [label="d b", color=red, shape=oval];
+ d_W [label="d W", color=red, shape=oval];
+
+ cost -> MSE_Grad [color=red];
+ d_cost -> MSE_Grad [color=red];
+ l -> MSE_Grad [color=red];
+ y -> MSE_Grad -> d_y [color=red];
+
+ x -> FC_Grad [color=red];
+ y -> FC_Grad [color=red];
+ d_y -> FC_Grad [color=red];
+ W -> FC_Grad -> d_W [color=red];
+ b -> FC_Grad -> d_b [color=red];
+
+ ////////// The optimizaiton part //////////
+
+ OPT_W [label="SGD", color=green, shape=box];
+ OPT_b [label="SGD", color=green, shape=box];
+
+ W -> OPT_W [color=green];
+ b -> OPT_b [color=green];
+ d_W -> OPT_W -> W [color=green];
+ d_b -> OPT_b -> b [color=green];
+
+ ////////// Groupings //////////
+
+ subgraph clusterMSE {
+ style=invis;
+ MSE;
+ MSE_Grad;
+ }
+
+ subgraph clusterFC {
+ style=invis;
+ FC;
+ FC_Grad;
+ }
+}
diff --git a/doc/fluid/images/graph_construction_example_all.png b/doc/fluid/images/graph_construction_example_all.png
new file mode 100644
index 0000000000..261611a572
Binary files /dev/null and b/doc/fluid/images/graph_construction_example_all.png differ
diff --git a/doc/fluid/images/graph_construction_example_forward_backward.png b/doc/fluid/images/graph_construction_example_forward_backward.png
new file mode 100644
index 0000000000..4c69687f4a
Binary files /dev/null and b/doc/fluid/images/graph_construction_example_forward_backward.png differ
diff --git a/doc/fluid/images/graph_construction_example_forward_only.png b/doc/fluid/images/graph_construction_example_forward_only.png
new file mode 100644
index 0000000000..e668c16e0c
Binary files /dev/null and b/doc/fluid/images/graph_construction_example_forward_only.png differ
diff --git a/doc/fluid/images/l1_regularization.png b/doc/fluid/images/l1_regularization.png
new file mode 100644
index 0000000000..e1b9c7a44f
Binary files /dev/null and b/doc/fluid/images/l1_regularization.png differ
diff --git a/doc/fluid/images/l2_regularization.png b/doc/fluid/images/l2_regularization.png
new file mode 100644
index 0000000000..d5c2fcbc2c
Binary files /dev/null and b/doc/fluid/images/l2_regularization.png differ
diff --git a/doc/fluid/images/local-graph.graffle b/doc/fluid/images/local-graph.graffle
new file mode 100644
index 0000000000..19e509bd9a
Binary files /dev/null and b/doc/fluid/images/local-graph.graffle differ
diff --git a/doc/fluid/images/local-graph.png b/doc/fluid/images/local-graph.png
new file mode 100644
index 0000000000..ada51200f7
Binary files /dev/null and b/doc/fluid/images/local-graph.png differ
diff --git a/doc/fluid/images/local_architecture.graffle b/doc/fluid/images/local_architecture.graffle
new file mode 100644
index 0000000000..49fcc663eb
Binary files /dev/null and b/doc/fluid/images/local_architecture.graffle differ
diff --git a/doc/fluid/images/local_architecture.png b/doc/fluid/images/local_architecture.png
new file mode 100644
index 0000000000..14adc9fd72
Binary files /dev/null and b/doc/fluid/images/local_architecture.png differ
diff --git a/doc/fluid/images/lookup_table.png b/doc/fluid/images/lookup_table.png
new file mode 100644
index 0000000000..72dfe3547f
Binary files /dev/null and b/doc/fluid/images/lookup_table.png differ
diff --git a/doc/fluid/images/lookup_table_training.png b/doc/fluid/images/lookup_table_training.png
new file mode 100644
index 0000000000..cc7cc4aeb3
Binary files /dev/null and b/doc/fluid/images/lookup_table_training.png differ
diff --git a/doc/fluid/images/loss_equation.png b/doc/fluid/images/loss_equation.png
new file mode 100644
index 0000000000..14212ec8d3
Binary files /dev/null and b/doc/fluid/images/loss_equation.png differ
diff --git a/doc/fluid/images/multi-threads.graffle b/doc/fluid/images/multi-threads.graffle
new file mode 100644
index 0000000000..e71173715f
Binary files /dev/null and b/doc/fluid/images/multi-threads.graffle differ
diff --git a/doc/fluid/images/multi-threads@3x.png b/doc/fluid/images/multi-threads@3x.png
new file mode 100644
index 0000000000..e40a869987
Binary files /dev/null and b/doc/fluid/images/multi-threads@3x.png differ
diff --git a/doc/fluid/images/multigpu_allreduce.graffle b/doc/fluid/images/multigpu_allreduce.graffle
new file mode 100644
index 0000000000..cb5bc420ce
Binary files /dev/null and b/doc/fluid/images/multigpu_allreduce.graffle differ
diff --git a/doc/fluid/images/multigpu_allreduce.png b/doc/fluid/images/multigpu_allreduce.png
new file mode 100644
index 0000000000..87a1b3e8f6
Binary files /dev/null and b/doc/fluid/images/multigpu_allreduce.png differ
diff --git a/doc/fluid/images/multigpu_before_convert.graffle b/doc/fluid/images/multigpu_before_convert.graffle
new file mode 100644
index 0000000000..6c35ab1b21
Binary files /dev/null and b/doc/fluid/images/multigpu_before_convert.graffle differ
diff --git a/doc/fluid/images/multigpu_before_convert.png b/doc/fluid/images/multigpu_before_convert.png
new file mode 100644
index 0000000000..9c8f771116
Binary files /dev/null and b/doc/fluid/images/multigpu_before_convert.png differ
diff --git a/doc/fluid/images/multiple_reader.png b/doc/fluid/images/multiple_reader.png
new file mode 100644
index 0000000000..b22126b31d
Binary files /dev/null and b/doc/fluid/images/multiple_reader.png differ
diff --git a/doc/fluid/images/paddle-compile.graffle b/doc/fluid/images/paddle-compile.graffle
new file mode 100644
index 0000000000..a6348cc3db
Binary files /dev/null and b/doc/fluid/images/paddle-compile.graffle differ
diff --git a/doc/fluid/images/paddle-compile.png b/doc/fluid/images/paddle-compile.png
new file mode 100644
index 0000000000..e0f13d551a
Binary files /dev/null and b/doc/fluid/images/paddle-compile.png differ
diff --git a/doc/fluid/images/pprof_1.png b/doc/fluid/images/pprof_1.png
new file mode 100644
index 0000000000..8e9edbf377
Binary files /dev/null and b/doc/fluid/images/pprof_1.png differ
diff --git a/doc/fluid/images/pprof_2.png b/doc/fluid/images/pprof_2.png
new file mode 100644
index 0000000000..172ba20399
Binary files /dev/null and b/doc/fluid/images/pprof_2.png differ
diff --git a/doc/fluid/images/profiler.png b/doc/fluid/images/profiler.png
new file mode 100644
index 0000000000..d57b71ca88
Binary files /dev/null and b/doc/fluid/images/profiler.png differ
diff --git a/doc/fluid/images/readers.png b/doc/fluid/images/readers.png
new file mode 100644
index 0000000000..fd59168ce1
Binary files /dev/null and b/doc/fluid/images/readers.png differ
diff --git a/doc/fluid/images/remote_executor.graffle b/doc/fluid/images/remote_executor.graffle
new file mode 100644
index 0000000000..41b2067311
Binary files /dev/null and b/doc/fluid/images/remote_executor.graffle differ
diff --git a/doc/fluid/images/remote_executor.png b/doc/fluid/images/remote_executor.png
new file mode 100644
index 0000000000..744e2fb2e0
Binary files /dev/null and b/doc/fluid/images/remote_executor.png differ
diff --git a/doc/fluid/images/rnn.dot b/doc/fluid/images/rnn.dot
new file mode 100644
index 0000000000..c1141cd9c9
--- /dev/null
+++ b/doc/fluid/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/fluid/images/rnn.jpg b/doc/fluid/images/rnn.jpg
new file mode 100644
index 0000000000..9867e404cf
Binary files /dev/null and b/doc/fluid/images/rnn.jpg differ
diff --git a/doc/fluid/images/rnn.png b/doc/fluid/images/rnn.png
new file mode 100644
index 0000000000..e139e373fe
Binary files /dev/null and b/doc/fluid/images/rnn.png differ
diff --git a/doc/fluid/images/rnn_2level_data.dot b/doc/fluid/images/rnn_2level_data.dot
new file mode 100644
index 0000000000..1d85ae2617
--- /dev/null
+++ b/doc/fluid/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/fluid/images/rnn_2level_data.png b/doc/fluid/images/rnn_2level_data.png
new file mode 100644
index 0000000000..4be81b2430
Binary files /dev/null and b/doc/fluid/images/rnn_2level_data.png differ
diff --git a/doc/fluid/images/single-thread@3x.png b/doc/fluid/images/single-thread@3x.png
new file mode 100644
index 0000000000..4083aebfdd
Binary files /dev/null and b/doc/fluid/images/single-thread@3x.png differ
diff --git a/doc/fluid/images/sparse_update.graffle b/doc/fluid/images/sparse_update.graffle
new file mode 100644
index 0000000000..08d689a58f
Binary files /dev/null and b/doc/fluid/images/sparse_update.graffle differ
diff --git a/doc/fluid/images/sparse_update.png b/doc/fluid/images/sparse_update.png
new file mode 100644
index 0000000000..8c872e6ac4
Binary files /dev/null and b/doc/fluid/images/sparse_update.png differ
diff --git a/doc/fluid/images/test.dot b/doc/fluid/images/test.dot
new file mode 100644
index 0000000000..62c69b8fc8
--- /dev/null
+++ b/doc/fluid/images/test.dot
@@ -0,0 +1,35 @@
+
+digraph Test {
+ z -> generator -> G_img;
+ G_img -> discriminator -> D_f -> d_loss_f;
+ label0 -> d_loss_f -> d_loss;
+
+ img -> discriminator -> D_t -> d_loss_t;
+ label1 -> d_loss_t -> d_loss;
+
+ d_loss -> d_loss_t[color=red, style=dashed];
+ d_loss -> d_loss_f[color=red, style=dashed];
+ d_loss_t -> D_t[color=red, style=dashed];
+ d_loss_f -> D_f[color=red, style=dashed];
+ D_t -> discriminator[color=red, style=dashed];
+ D_f -> discriminator[color=red, style=dashed];
+
+ D_f -> g_loss;
+ label2 -> g_loss;
+
+ g_loss -> D_f[color=green, style=dashed];
+ D_f -> discriminator[color=green, style=dashed];
+ discriminator -> G_img[color=green, style=dashed];
+ G_img -> generator[color=green, style=dashed];
+
+ discriminator [color=red, shape=box];
+ generator [color=green, shape=box];
+ z [shape=diamond];
+ img [shape=diamond];
+ label0 [shape=diamond];
+ label1 [shape=diamond];
+ label2 [shape=diamond];
+
+ d_loss [color=red];
+ g_loss [color=green];
+}
diff --git a/doc/fluid/images/test.dot.png b/doc/fluid/images/test.dot.png
new file mode 100644
index 0000000000..4e121a40b9
Binary files /dev/null and b/doc/fluid/images/test.dot.png differ
diff --git a/doc/fluid/images/theta_star.gif b/doc/fluid/images/theta_star.gif
new file mode 100644
index 0000000000..dd24d33e12
Binary files /dev/null and b/doc/fluid/images/theta_star.gif differ
diff --git a/doc/fluid/images/timeline.jpeg b/doc/fluid/images/timeline.jpeg
new file mode 100644
index 0000000000..38ec3f80c9
Binary files /dev/null and b/doc/fluid/images/timeline.jpeg differ
diff --git a/doc/fluid/images/tracing.jpeg b/doc/fluid/images/tracing.jpeg
new file mode 100644
index 0000000000..3a49fc4f8a
Binary files /dev/null and b/doc/fluid/images/tracing.jpeg differ
diff --git a/doc/v2/howto/rnn/recurrent_group_en.md b/doc/v2/howto/rnn/recurrent_group_en.md
index d264b0a9f8..de6b60f29e 100644
--- a/doc/v2/howto/rnn/recurrent_group_en.md
+++ b/doc/v2/howto/rnn/recurrent_group_en.md
@@ -1,3 +1,96 @@
# Recurrent Group Tutorial
-TBD
+## Overview
+
+Sequential data is common in natural language processing.
+
+A sentence is a sequence of words and many sentences form a paragraph further. Therefore, a paragraph can be viewed as a nested sequence with two level, where each element of the sequence is another sequence. That is to say, sequential data could be recursive. An example of two-level recursive sequential data is that an article is composed of a sequence of sentences, and each sentence a sequence of words.
+
+PaddlePaddle and PaddlePaddle v2 support two-level recursive sequential data. The two-level sequence is a very flexible data, which helps us to better describe more complex language data such as discribing paragraphs and several rounds of dialogues. Based on two-level sequence input, we can design and build a flexible, hierarchical RNN model that encodes input data from the word and sentence level. For the support of arbitrary levels, please refer to PaddlePaddle Fluid.
+
+In PaddlePaddle, `recurrent_group` is an arbitrarily complex RNN unit. The user only needs to define the calculation that the RNN will complete in one time step. PaddlePaddle is responsible for the propagation of information and error in time series.
+
+Furthermore, `recurrent_group` can also be extended to handle two-level sequence. By defining two nested `recurrent_group` operations at the clause level and the word level respectively, a hierarchical and complex RNN is finally achieved.
+
+Currently, in the PaddlePaddle, there are `recurrent_group` and some Layers that can process bidirectional sequences. For details, refer to the document: Layers for supporting double-layer sequences as input.
+
+## Related Concepts
+
+### Basic Principle
+`recurrent_group` is an arbitrarily complex RNN unit supported by PaddlePaddle. The user only needs to focus on the calculations that the RNN is designed to complete within a single time step. The PaddlePaddle is responsible for completing the propagation of information and gradients over time.
+
+In PaddlePaddle, a simple call to `recurrent_group` is as follows:
+
+``` python
+recurrent_group(step, input, reverse)
+```
+- step: A callable function that defines the calculations completed by the RNN unit within a time step
+- input: The input must be a single-layer sequence or a double-layer sequence
+- reverse: Whether to process the input sequence in reverse order
+
+The core of using `recurrent_group` is to design the logic of the step function. The step function can be freely combined with various layers supported by PaddlePaddle to complete arbitrary arithmetic logic. The input of `recurrent_group` (input) becomes the input of the step function. Since the step function only focuses on the calculation within one time step of RNN, here `recurrent_group` completes the splitting of the original input data for us.
+
+### Input
+The input sequence processed by `recurrent_group` is mainly divided into the following three types:
+
+- **Input Data**: When putting a two-level sequence into `recurrent_group`, it will be disassembled into a single-level sequence. When putting a single-level sequence into `recurrent_group`, it will be disassembled into a non-sequence and then passed to the step function. This process is completely transparent to the user. There are two possible types: 1) User input via data_layer; 2) Output from other layers.
+
+- **Read-only Memory Input**: `StaticInput` defines a read-only Memory. The input specified by `StaticInput` will not be disassembled by `recurrent_group`, and each time step of the `recurrent_group` loop will always be able to reference all inputs. It may be a non-sequence or a single-layer sequence.
+
+- **Input of Sequence Generation Task**: `GeneratedInput` is only used to specify input data in a sequence generation task.
+
+### Input Example
+
+Sequence generation tasks mostly follow the encoder-decoer architecture. The encoder and decoder can be arbitrary neural network units capable of processing sequences and RNN is the most popular choice.
+
+Given the encoder output and the current word, the decoder predicts the next most likely word each time. In this structure, the decoder accepts two inputs:
+
+- Target sequence to be generated: a input of the decoder and the basis of the decoder loop. `recurrent_group` will disassemble this input type.
+
+- Encoder output, an non-sequencce or single-sequence: a unbounded memory. Each time step in the decoder loop will reference the entire result and should not be disassembled. This type of input must be specified via `StaticInput`. For more discussion on Unbounded Memory, please refer to the paper [Neural Turning Machine](https://arxiv.org/abs/1410.5401).
+
+In a sequence generation task, the decoder RNN always refers to the word vector of the word predicted at the previous moment as the current time input. `GeneratedInput` will automate this process.
+
+### Output
+The `step` function must return the output of one or more Layers. The output of this Layer will be the final output of the entire `recurrent_group`. In the output process, `recurrent_group` will concatenate the output of each time step, which is also transparent to the user.
+
+### Memory
+Memory can only be defined and used in `recurrent_group`. Memory cannot exist independently and must point to a layer defined by PaddlePaddle. Memory is referenced to get a momentary output from this layer, so memory can be interpreted as a delay operation.
+
+The user can explicitly specify the output of a layer to initialize the memory. When not specified, memory is initialized to 0 by default.
+
+## Sequence-level RNN Introduction
+
+`recurrent_group` helps us to split the input sequence, merge the output, and loop through the sequence of computational logic.
+
+Using this feature, the two nested `recurrent_group` can handle the nested two-level sequences, implementing sequence-level RNN structures at both the word and sentence levels.
+
+- Word-level RNN: each state corresponds to a word.
+- Sequence-level RNN: a sequence-layer RNN consists of multiple word-layer RNNs. Each word-layer RNN (ie, each state of a sequence-layer RNN) has a subsequence.
+
+For convenience of description, the following takes the NLP task as an example. A paragraph containing a subsequence is defined as a two-level sequence, and a sentence containing a word is defined as a single-layer sequence. Then, the zero-level sequence is a word.
+
+## Usage of Sequence-level RNN
+
+### Usage of Training Process
+Using `recurrent_group` requires the following conventions:
+
+- **Single-input Single-output**: Both input and output are single layer sequences.
+ - If there are multiple inputs, the number of words in different input sequences must be exactly equal.
+ - A single-layer sequence is output, and the number of words in the output sequence is the same as the input sequence.
+ - memory: define memory to point to a layer in the step function, get a moment output from this layer by referencing memory to form a recurrent connection. The is_seq parameter of memory must be false. If memory is not defined, the operations within each time step are independent.
+ - boot_layer: the initial state of memory, set 0 by default. is_seq in memory must be false.
+
+- **Double-input Double-output**: Both input and output are two-level sequence.
+ - If there are multiple input sequences, the number of subsequence contained in different inputs must be strictly equal, but the number of words in the subsequence may not be equal.
+ - output a two-level sequence. The number of subsequence and the number of words are the same as the specified input sequence and the first input is default.
+ - memory: defining memory in the step function, pointing to a layer, by referring to the memory to get the output of this layer at a time, forming a recurrent connection. The memory defined in the outer `recurrent_group` step function can record the state of the previous subsequence, either as a single-level sequence (only as read-only memory) or as a word. If memory is not defined, the operations between subsequence are independent.
+ - boot_layer: the initial state of memory. It is either a single-level sequence (only as read-only memory) or a vector. The default is not set, that is, the initial state is 0.
+
+- **Double-input Single-output**: not support for now, and output the error with "In hierachical RNN, all out links should be from sequences now".
+
+### Usage of Generation Process
+Using `beam_search` need follow those conventions:
+
+- Word-level RNN: generate the next word from a word.
+- Sequence-level RNN: the single-layer RNN generated subsequence is concatenated into a new double-layer sequence. Semantically, there is no case where a subsequence generates the next subseq directly.
diff --git a/paddle/cuda/include/hl_cnn.h b/paddle/cuda/include/hl_cnn.h
index 63ec515647..b790fa39fe 100644
--- a/paddle/cuda/include/hl_cnn.h
+++ b/paddle/cuda/include/hl_cnn.h
@@ -370,4 +370,48 @@ extern void hl_maxout_backward(real* inGrad,
size_t featLen,
size_t groups);
+/**
+ * @brief Upsample forward.
+ * @param[in] inputData input data.
+ * @param[out] maskData the mask data from MaxPoolWithMaskLayer.
+ * @param[out] batchSize the batch size of the input.
+ * @param[in] imgSizeH image height.
+ * @param[in] imgSizeW image width.
+ * @param[in] channels the input channels.
+ * @param[in] outputH the output height.
+ * @param[in] outputW the output widht.
+ * @param[out] outputData output data.
+ */
+extern void hl_upsample_forward(real* inputData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* outputData);
+
+/**
+ * @brief Upsample backward.
+ * @param[in] outputGradData the output grad data.
+ * @param[out] maskData the mask data from MaxPoolWithMaskLayer.
+ * @param[out] batchSize the batch size of the input.
+ * @param[in] imgSizeH image height.
+ * @param[in] imgSizeW image width.
+ * @param[in] channels the input channels.
+ * @param[in] outputH the output height.
+ * @param[in] outputW the output widht.
+ * @param[out] inputGradData the input grad data.
+ */
+extern void hl_upsample_backward(real* outputGradData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* inputGradData);
+
#endif // HL_CNN_H_
diff --git a/paddle/cuda/include/stub/hl_cnn_stub.h b/paddle/cuda/include/stub/hl_cnn_stub.h
index c39bd3228d..997eed62e0 100644
--- a/paddle/cuda/include/stub/hl_cnn_stub.h
+++ b/paddle/cuda/include/stub/hl_cnn_stub.h
@@ -224,4 +224,24 @@ inline void hl_maxout_backward(real* inGrad,
size_t featLen,
size_t group) {}
+inline void hl_upsample_forward(real* inputData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* outputData) {}
+
+inline void hl_upsample_backward(real* outputGradData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* inputGradData) {}
+
#endif // HL_CNN_STUB_H_
diff --git a/paddle/cuda/src/hl_cuda_cnn.cu b/paddle/cuda/src/hl_cuda_cnn.cu
index a4459243e8..bac743a293 100644
--- a/paddle/cuda/src/hl_cuda_cnn.cu
+++ b/paddle/cuda/src/hl_cuda_cnn.cu
@@ -1028,3 +1028,79 @@ void hl_maxout_backward(real* inGrad,
num_kernels, inGrad, outGrad, idData, size, featLen, groups);
CHECK_SYNC("hl_maxout_backward failed");
}
+
+__global__ void upsampleForwardCompute(real* input_data,
+ real* mask_data,
+ size_t nthreads,
+ size_t in_h,
+ size_t in_w,
+ size_t out_h,
+ size_t out_w,
+ real* output_data) {
+ int index = blockIdx.x * blockDim.x + threadIdx.x;
+ if (index < nthreads) {
+ int offset = index / (in_w * in_h) * out_h * out_w;
+ int upsample_idx = static_cast(mask_data[index]);
+ output_data[offset + upsample_idx] = input_data[index];
+ }
+}
+
+__global__ void upsampleBackwardCompute(real* out_grad,
+ real* mask_data,
+ size_t nthreads,
+ size_t in_h,
+ size_t in_w,
+ size_t out_h,
+ size_t out_w,
+ real* input_grad) {
+ int index = blockIdx.x * blockDim.x + threadIdx.x;
+ if (index < nthreads) {
+ int offset = index / (in_w * in_h) * out_h * out_w;
+ int upsample_idx = static_cast(mask_data[index]);
+ input_grad[index] = out_grad[offset + upsample_idx];
+ }
+}
+
+void hl_upsample_forward(real* inputData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* outputData) {
+ int num_kernels = batchSize * imgSizeH * imgSizeW * channels;
+ int blocks = (num_kernels + 1024 - 1) / 1024;
+ upsampleForwardCompute<<>>(inputData,
+ maskData,
+ num_kernels,
+ imgSizeH,
+ imgSizeW,
+ outputH,
+ outputW,
+ outputData);
+ CHECK_SYNC("hl_upsample_forward failed");
+}
+
+void hl_upsample_backward(real* outputGradData,
+ real* maskData,
+ size_t batchSize,
+ size_t imgSizeH,
+ size_t imgSizeW,
+ size_t channels,
+ size_t outputH,
+ size_t outputW,
+ real* inputGradData) {
+ int num_kernels = batchSize * imgSizeH * imgSizeW * channels;
+ int blocks = (num_kernels + 1024 - 1) / 1024;
+ upsampleBackwardCompute<<>>(outputGradData,
+ maskData,
+ num_kernels,
+ imgSizeH,
+ imgSizeW,
+ outputH,
+ outputW,
+ inputGradData);
+ CHECK_SYNC("hl_upsample_backward failed");
+}
diff --git a/paddle/fluid/framework/details/multi_devices_graph_builder.cc b/paddle/fluid/framework/details/multi_devices_graph_builder.cc
index a1b913a863..128a5344fb 100644
--- a/paddle/fluid/framework/details/multi_devices_graph_builder.cc
+++ b/paddle/fluid/framework/details/multi_devices_graph_builder.cc
@@ -21,6 +21,9 @@
#include "paddle/fluid/framework/details/nccl_all_reduce_op_handle.h"
#endif
+#include
+#include
+
namespace paddle {
namespace framework {
namespace details {
@@ -55,6 +58,7 @@ std::unique_ptr MultiDevSSAGraphBuilder::Build(
const ProgramDesc &program) const {
auto graph = new SSAGraph();
SSAGraph &result = *graph;
+ std::unordered_set og_has_been_broadcast;
result.vars_.resize(places_.size());
bool is_forwarding = true;
@@ -122,9 +126,15 @@ std::unique_ptr MultiDevSSAGraphBuilder::Build(
if (!is_forwarding) {
auto var_names = op->OutputArgumentNames();
+ // Currently, we assume that once gradient is generated, it can be
+ // broadcast, and each gradient is only broadcast once. But there are no
+ // other cases, for example, we need to adjust the gradient according to
+ // the input when we get the gradient, which is not considered at present.
for (auto &og : var_names) {
- if (grad_names_.count(og) != 0) { // is param grad
- // Insert NCCL AllReduce Op
+ if (grad_names_.count(og) != 0 &&
+ og_has_been_broadcast.count(og) == 0) { // is param grad
+ // Insert NCCL AllReduce Op
+ og_has_been_broadcast.insert(og);
#ifdef PADDLE_WITH_CUDA
result.ops_.emplace_back(
new NCCLAllReduceOpHandle(local_scopes_, places_, *nccl_ctxs_));
@@ -161,6 +171,11 @@ std::unique_ptr MultiDevSSAGraphBuilder::Build(
*/
PolishGraphToSupportDataHazards(&result);
+ /*
+ * Only variables should be the leaves of graph.
+ */
+ AddOutputToLeafOps(&result);
+
if (VLOG_IS_ON(10)) {
std::ostringstream sout;
PrintGraphviz(*graph, sout);
diff --git a/paddle/fluid/framework/details/nccl_all_reduce_op_handle.cc b/paddle/fluid/framework/details/nccl_all_reduce_op_handle.cc
index 5ddf331cfc..55b5f11358 100644
--- a/paddle/fluid/framework/details/nccl_all_reduce_op_handle.cc
+++ b/paddle/fluid/framework/details/nccl_all_reduce_op_handle.cc
@@ -76,7 +76,7 @@ void NCCLAllReduceOpHandle::RunImpl() {
}
}
-std::string NCCLAllReduceOpHandle::Name() const { return "NCCL AllReduce"; }
+std::string NCCLAllReduceOpHandle::Name() const { return "nccl_all_reduce"; }
} // namespace details
} // namespace framework
} // namespace paddle
diff --git a/paddle/fluid/framework/details/nccl_all_reduce_op_handle.h b/paddle/fluid/framework/details/nccl_all_reduce_op_handle.h
index 045070bb6a..ad14a3c5cb 100644
--- a/paddle/fluid/framework/details/nccl_all_reduce_op_handle.h
+++ b/paddle/fluid/framework/details/nccl_all_reduce_op_handle.h
@@ -14,6 +14,9 @@
#pragma once
+#include
+#include
+
#include "paddle/fluid/framework/details/op_handle_base.h"
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/scope.h"
@@ -34,6 +37,10 @@ struct NCCLAllReduceOpHandle : public OpHandleBase {
std::string Name() const override;
+ // Delay and buffer nccl_all_reduce together can significantly increase
+ // performance. Disable this feature by returning false.
+ bool IsMultiDeviceTransfer() override { return true; };
+
protected:
void RunImpl() override;
};
diff --git a/paddle/fluid/framework/details/op_handle_base.h b/paddle/fluid/framework/details/op_handle_base.h
index 71672fd24c..d7a541ac4b 100644
--- a/paddle/fluid/framework/details/op_handle_base.h
+++ b/paddle/fluid/framework/details/op_handle_base.h
@@ -13,6 +13,8 @@
// limitations under the License.
#pragma once
+#include
+#include
#include "paddle/fluid/framework/details/var_handle.h"
#include "paddle/fluid/platform/device_context.h"
@@ -53,6 +55,10 @@ class OpHandleBase {
void AddOutput(VarHandleBase *out);
+ // If the Op involves data transfer of multiple devices that
+ // will likely block other computations.
+ virtual bool IsMultiDeviceTransfer() { return false; }
+
protected:
virtual void RunImpl() = 0;
};
diff --git a/paddle/fluid/framework/details/ssa_graph_builder.cc b/paddle/fluid/framework/details/ssa_graph_builder.cc
index 361ba6d397..0a4febd22f 100644
--- a/paddle/fluid/framework/details/ssa_graph_builder.cc
+++ b/paddle/fluid/framework/details/ssa_graph_builder.cc
@@ -136,6 +136,17 @@ void SSAGraphBuilder::PrintGraphviz(const SSAGraph &graph, std::ostream &sout) {
sout << "}\n";
}
+
+void SSAGraphBuilder::AddOutputToLeafOps(SSAGraph *graph) {
+ for (auto &op : graph->ops_) {
+ if (!op->outputs_.empty()) {
+ continue;
+ }
+ auto *dummy_leaf = new DummyVarHandle();
+ graph->dep_vars_.emplace(dummy_leaf);
+ op->AddOutput(dummy_leaf);
+ }
+}
} // namespace details
} // namespace framework
} // namespace paddle
diff --git a/paddle/fluid/framework/details/ssa_graph_builder.h b/paddle/fluid/framework/details/ssa_graph_builder.h
index bf20e7164a..be1f0460e4 100644
--- a/paddle/fluid/framework/details/ssa_graph_builder.h
+++ b/paddle/fluid/framework/details/ssa_graph_builder.h
@@ -14,13 +14,13 @@
#pragma once
+#include
+#include
+
#include "paddle/fluid/framework/details/ssa_graph.h"
#include "paddle/fluid/framework/program_desc.h"
#include "paddle/fluid/platform/place.h"
-#include
-#include
-
namespace paddle {
namespace framework {
namespace details {
@@ -52,6 +52,8 @@ class SSAGraphBuilder {
const std::string &each_var_name,
const platform::Place &place, size_t place_offset);
+ static void AddOutputToLeafOps(SSAGraph *graph);
+
static void PrintGraphviz(const SSAGraph &graph, std::ostream &sout);
};
} // namespace details
diff --git a/paddle/fluid/framework/details/threaded_ssa_graph_executor.cc b/paddle/fluid/framework/details/threaded_ssa_graph_executor.cc
index 3f8655147b..596e573186 100644
--- a/paddle/fluid/framework/details/threaded_ssa_graph_executor.cc
+++ b/paddle/fluid/framework/details/threaded_ssa_graph_executor.cc
@@ -23,22 +23,36 @@ ThreadedSSAGraphExecutor::ThreadedSSAGraphExecutor(
size_t num_threads, bool use_event,
const std::vector &local_scopes,
const std::vector &places,
- std::unique_ptr &&graph)
+ std::unique_ptr &&graph, bool allow_op_delay)
: SSAGraphExecutor(std::move(graph)),
pool_(num_threads >= 2 ? new ::ThreadPool(num_threads) : nullptr),
local_scopes_(local_scopes),
places_(places),
fetch_ctxs_(places),
- use_event_(use_event) {}
+ use_event_(use_event),
+ running_ops_(0),
+ allow_op_delay_(allow_op_delay) {}
+
+void ThreadedSSAGraphExecutor::RunDelayedOps(
+ const std::unordered_set &delayed_ops) {
+ for (auto op : delayed_ops) {
+ op->Run(use_event_);
+ }
+}
FeedFetchList ThreadedSSAGraphExecutor::Run(
const std::vector &fetch_tensors) {
std::unordered_map pending_ops;
std::unordered_set pending_vars;
-
BlockingQueue ready_vars;
-
std::unordered_set ready_ops;
+ // For ops (e.g. nccl_all_reduce) that need to coordinate multiple
+ // streams from multiple GPUs, it's faster to buffer them and schedule
+ // together since we currently cannot overlap computation and memcpy streams.
+ // Should revisit it if overlapping is available.
+ std::unordered_set delayed_ops;
+ std::unordered_set blocked_by_delayed_ops;
+ std::unordered_set delayed_vars;
auto InsertPendingVar = [&pending_vars, &ready_vars](VarHandleBase &var) {
pending_vars.insert(&var);
@@ -73,7 +87,6 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
// Step 2. Insert FetchOps
std::vector> fetch_ops;
- std::vector dummy_vars;
FeedFetchList fetch_data(fetch_tensors.size());
std::unordered_map> fetched_vars;
@@ -87,13 +100,13 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
}
}
+ std::unordered_set> fetch_dependencies;
for (size_t i = 0; i < fetch_tensors.size(); ++i) {
auto &var_name = fetch_tensors[i];
auto &vars = fetched_vars.at(var_name);
auto *op = new FetchOpHandle(&fetch_data, i, &local_scopes_);
fetch_ops.emplace_back(op);
- // FIXME: Use new device context
for (auto &p : places_) {
op->dev_ctxes_[p] = fetch_ctxs_.Get(p);
}
@@ -101,12 +114,24 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
for (auto *var : vars) {
op->AddInput(var);
}
+
+ auto *fetch_dummy = new DummyVarHandle();
+ op->AddOutput(fetch_dummy);
+ fetch_dependencies.emplace(fetch_dummy);
+ InsertPendingVar(*fetch_dummy);
InsertPendingOp(*op);
}
auto run_all_ready_ops = [&] {
for (auto *op : ready_ops) {
- RunOp(ready_vars, op);
+ if (op->IsMultiDeviceTransfer() && allow_op_delay_) {
+ delayed_ops.insert(op);
+ delayed_vars.insert(op->outputs_.begin(), op->outputs_.end());
+ ready_vars.Extend(op->outputs_);
+ continue;
+ }
+ running_ops_++;
+ RunOp(&ready_vars, op);
}
ready_ops.clear();
};
@@ -118,13 +143,13 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
}
// Step 3. Execution
- while (!pending_vars.empty()) {
+ while (!pending_vars.empty() || !ready_ops.empty() || !delayed_ops.empty()) {
// 1. Run All Ready ops
run_all_ready_ops();
// 2. Find ready variable
bool timeout;
- auto cur_ready_vars = ready_vars.PopAll(1000, &timeout);
+ auto cur_ready_vars = ready_vars.PopAll(1, &timeout);
if (timeout) {
if (exception_) {
@@ -141,13 +166,29 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
auto &deps = pending_ops[op];
--deps;
if (deps == 0) {
- ready_ops.insert(op);
+ if (delayed_vars.find(ready_var) != delayed_vars.end()) {
+ blocked_by_delayed_ops.insert(op);
+ } else {
+ ready_ops.insert(op);
+ }
}
}
}
+ // When there are no other ops to schedule, schedule buffered delayed
+ // ops and unblock other ops.
+ if (ready_ops.empty() && !delayed_ops.empty() && running_ops_ == 0) {
+ RunDelayedOps(delayed_ops);
+ delayed_ops.clear();
+ for (auto *op : blocked_by_delayed_ops) {
+ ready_ops.insert(op);
+ }
+ blocked_by_delayed_ops.clear();
+ }
// Keep loop until all vars are ready.
}
-
+ PADDLE_ENFORCE(ready_ops.empty());
+ PADDLE_ENFORCE(delayed_ops.empty());
+ PADDLE_ENFORCE(blocked_by_delayed_ops.empty());
++computation_count_;
auto sync_computation = [&] {
@@ -182,12 +223,13 @@ FeedFetchList ThreadedSSAGraphExecutor::Run(
}
void ThreadedSSAGraphExecutor::RunOp(
- BlockingQueue &ready_var_q, details::OpHandleBase *op) {
- auto op_run = [&ready_var_q, op, this] {
+ BlockingQueue *ready_var_q, details::OpHandleBase *op) {
+ auto op_run = [ready_var_q, op, this] {
try {
VLOG(10) << op->Name() << " : " << op->DebugString();
op->Run(use_event_);
- ready_var_q.Extend(op->outputs_);
+ running_ops_--;
+ ready_var_q->Extend(op->outputs_);
} catch (platform::EnforceNotMet ex) {
exception_.reset(new platform::EnforceNotMet(ex));
} catch (...) {
diff --git a/paddle/fluid/framework/details/threaded_ssa_graph_executor.h b/paddle/fluid/framework/details/threaded_ssa_graph_executor.h
index 2ea57ac8f9..79cfc26b46 100644
--- a/paddle/fluid/framework/details/threaded_ssa_graph_executor.h
+++ b/paddle/fluid/framework/details/threaded_ssa_graph_executor.h
@@ -14,7 +14,12 @@
#pragma once
-#include
+#include
+#include
+#include
+#include
+#include
+
#include
#include "ThreadPool.h" // ThreadPool in thrird party
#include "paddle/fluid/framework/details/ssa_graph_executor.h"
@@ -70,7 +75,8 @@ class ThreadedSSAGraphExecutor : public SSAGraphExecutor {
ThreadedSSAGraphExecutor(size_t num_threads, bool use_event,
const std::vector &local_scopes,
const std::vector &places,
- std::unique_ptr &&graph);
+ std::unique_ptr &&graph,
+ bool allow_op_delay);
// Run a SSAGraph by a thread pool
// Use topological sort algorithm
@@ -79,9 +85,11 @@ class ThreadedSSAGraphExecutor : public SSAGraphExecutor {
~ThreadedSSAGraphExecutor() {}
private:
- void RunOp(BlockingQueue &ready_var_q,
+ void RunOp(BlockingQueue *ready_var_q,
details::OpHandleBase *op);
+ void RunDelayedOps(const std::unordered_set &delayed_ops);
+
private:
std::unique_ptr<::ThreadPool> pool_;
std::vector local_scopes_;
@@ -89,6 +97,8 @@ class ThreadedSSAGraphExecutor : public SSAGraphExecutor {
platform::DeviceContextPool fetch_ctxs_;
const bool use_event_;
std::unique_ptr exception_;
+ std::atomic running_ops_;
+ bool allow_op_delay_;
size_t computation_count_{0};
size_t max_async_computation{100};
diff --git a/paddle/fluid/framework/executor.cc b/paddle/fluid/framework/executor.cc
index 009d0fbeb8..8a0ab118d0 100644
--- a/paddle/fluid/framework/executor.cc
+++ b/paddle/fluid/framework/executor.cc
@@ -265,6 +265,21 @@ std::unique_ptr Executor::Prepare(
return std::unique_ptr(ctx);
}
+std::vector> Executor::Prepare(
+ const ProgramDesc& program, const std::vector& block_ids) {
+ std::vector> result;
+ for (auto& bid : block_ids) {
+ auto* ctx = new ExecutorPrepareContext(program, bid);
+ PADDLE_ENFORCE_LT(static_cast(bid), program.Size());
+ auto& block = program.Block(bid);
+ for (auto& op_desc : block.AllOps()) {
+ ctx->ops_.push_back(OpRegistry::CreateOp(*op_desc));
+ }
+ result.push_back(std::shared_ptr(ctx));
+ }
+ return result;
+}
+
void Executor::RunPreparedContext(ExecutorPrepareContext* ctx, Scope* scope,
bool create_local_scope, bool create_vars) {
auto& block = ctx->prog_.Block(ctx->block_id_);
diff --git a/paddle/fluid/framework/executor.h b/paddle/fluid/framework/executor.h
index b0e64d5de0..462a56d49f 100644
--- a/paddle/fluid/framework/executor.h
+++ b/paddle/fluid/framework/executor.h
@@ -61,6 +61,9 @@ class Executor {
static std::unique_ptr Prepare(
const ProgramDesc& program, int block_id);
+ static std::vector> Prepare(
+ const ProgramDesc& program, const std::vector& block_ids);
+
void RunPreparedContext(ExecutorPrepareContext* ctx, Scope* scope,
bool create_local_scope = true,
bool create_vars = true);
diff --git a/paddle/fluid/framework/parallel_executor.cc b/paddle/fluid/framework/parallel_executor.cc
index 577eea92d2..1788514324 100644
--- a/paddle/fluid/framework/parallel_executor.cc
+++ b/paddle/fluid/framework/parallel_executor.cc
@@ -13,6 +13,7 @@ See the License for the specific language governing permissions and
limitations under the License. */
#include "paddle/fluid/framework/parallel_executor.h"
+#include "paddle/fluid/platform/profiler.h"
#include
#include
@@ -47,7 +48,7 @@ ParallelExecutor::ParallelExecutor(
const std::vector &places,
const std::unordered_set ¶ms,
const ProgramDesc &startup_program, const ProgramDesc &main_program,
- const std::string &loss_var_name, Scope *scope)
+ const std::string &loss_var_name, Scope *scope, bool allow_op_delay)
: member_(new ParallelExecutorPrivate(places)) {
member_->global_scope_ = scope;
@@ -82,8 +83,8 @@ ParallelExecutor::ParallelExecutor(
auto graph = builder.Build(main_program);
member_->executor_.reset(new details::ThreadedSSAGraphExecutor(
- num_threads, use_event, member_->local_scopes_, places,
- std::move(graph)));
+ num_threads, use_event, member_->local_scopes_, places, std::move(graph),
+ allow_op_delay));
// Step 3. Create vars in each scope;
for (auto *scope : member_->local_scopes_) {
@@ -151,6 +152,7 @@ void ParallelExecutor::BCastParamsToGPUs(
void ParallelExecutor::Run(const std::vector &fetch_tensors,
const std::string &fetched_var_name) {
+ platform::RecordBlock b(0);
auto fetch_data = member_->executor_->Run(fetch_tensors);
*member_->global_scope_->Var(fetched_var_name)->GetMutable() =
fetch_data;
diff --git a/paddle/fluid/framework/parallel_executor.h b/paddle/fluid/framework/parallel_executor.h
index 503efa2e44..964b476234 100644
--- a/paddle/fluid/framework/parallel_executor.h
+++ b/paddle/fluid/framework/parallel_executor.h
@@ -14,8 +14,9 @@ limitations under the License. */
#pragma once
-#include
+#include
#include
+#include
#include "paddle/fluid/framework/executor.h"
#include "paddle/fluid/framework/op_info.h"
#include "paddle/fluid/framework/program_desc.h"
@@ -37,7 +38,8 @@ class ParallelExecutor {
const std::unordered_set& params,
const ProgramDesc& startup_program,
const ProgramDesc& main_program,
- const std::string& loss_var_name, Scope* scope);
+ const std::string& loss_var_name, Scope* scope,
+ bool allow_op_delay);
void Run(const std::vector& fetch_tensors,
const std::string& fetched_var_name = "fetched_var");
diff --git a/paddle/fluid/framework/selected_rows.h b/paddle/fluid/framework/selected_rows.h
index c9c2c1bb72..9458d56a01 100644
--- a/paddle/fluid/framework/selected_rows.h
+++ b/paddle/fluid/framework/selected_rows.h
@@ -10,6 +10,9 @@ See the License for the specific language governing permissions and
limitations under the License. */
#pragma once
+
+#include
+
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/tensor.h"
@@ -52,7 +55,7 @@ class SelectedRows {
private:
// Notice: rows can be duplicate. We can have {0, 4, 7, 0, 5, 7, 9} here.
- // SelectedRows are simplely concated when adding together. Until a
+ // SelectedRows are simply concated when adding together. Until a
// SelectedRows add a Tensor, will the duplicate rows be handled.
Vector rows_;
std::unique_ptr value_{nullptr};
diff --git a/paddle/fluid/framework/tensor.h b/paddle/fluid/framework/tensor.h
index f7a6b5ba84..6f878541e6 100644
--- a/paddle/fluid/framework/tensor.h
+++ b/paddle/fluid/framework/tensor.h
@@ -45,11 +45,10 @@ class Tensor {
friend struct EigenVector;
public:
- Tensor() : offset_(0), is_pinned_(false) {}
+ Tensor() : offset_(0) {}
/*! Constructor with place should only be used in pybind. */
- explicit Tensor(const platform::Place& place)
- : offset_(0), is_pinned_(false) {
+ explicit Tensor(const platform::Place& place) : offset_(0) {
holder_->set_place(place);
}
@@ -70,12 +69,11 @@ class Tensor {
* @note If not exist, then allocation.
*/
template
- inline T* mutable_data(platform::Place place, bool is_pinned = false);
+ inline T* mutable_data(platform::Place place);
- inline void* mutable_data(platform::Place place, std::type_index type,
- bool is_pinned = false);
+ inline void* mutable_data(platform::Place place, std::type_index type);
- inline void* mutable_data(platform::Place place, bool is_pinned = false);
+ inline void* mutable_data(platform::Place place);
/**
* @brief Return a pointer to mutable memory block.
@@ -86,8 +84,7 @@ class Tensor {
* @note If not exist, then allocation.
*/
template
- inline T* mutable_data(DDim dims, platform::Place place,
- bool is_pinned = false);
+ inline T* mutable_data(DDim dims, platform::Place place);
/*! Return the dimensions of the memory block. */
inline const DDim& dims() const;
@@ -95,9 +92,6 @@ class Tensor {
/*! Return the numel of the memory block. */
inline int64_t numel() const;
- /*! Return the numel of the memory block. */
- inline bool isPinned() const;
-
/*! Resize the dimensions of the memory block. */
inline Tensor& Resize(const DDim& dims);
@@ -152,14 +146,12 @@ class Tensor {
template
struct PlaceholderImpl : public Placeholder {
- PlaceholderImpl(Place place, size_t size, std::type_index type,
- bool is_pinned = false)
- : ptr_(static_cast(memory::Alloc(place, size, is_pinned)),
- memory::PODDeleter(place, is_pinned)),
+ PlaceholderImpl(Place place, size_t size, std::type_index type)
+ : ptr_(static_cast(memory::Alloc(place, size)),
+ memory::PODDeleter(place)),
place_(place),
size_(size),
- type_(type),
- is_pinned_(is_pinned) {
+ type_(type) {
PADDLE_ENFORCE_NOT_NULL(ptr_, "Insufficient %s memory to allocation.",
(is_cpu_place(place_) ? "CPU" : "GPU"));
}
@@ -182,9 +174,6 @@ class Tensor {
/* the current type of memory */
std::type_index type_;
-
- /*! use pinned memory or not. */
- bool is_pinned_;
};
/*! holds the memory block if allocated. */
@@ -219,7 +208,6 @@ class Tensor {
* PlaceHolder::ptr_ and where the tensor data really begins.
*/
size_t offset_;
- bool is_pinned_;
};
inline void Tensor::switch_place(platform::Place new_place) {
diff --git a/paddle/fluid/framework/tensor_impl.h b/paddle/fluid/framework/tensor_impl.h
index 113814971e..7a48390440 100644
--- a/paddle/fluid/framework/tensor_impl.h
+++ b/paddle/fluid/framework/tensor_impl.h
@@ -101,21 +101,19 @@ inline T* Tensor::data() {
}
template
-inline T* Tensor::mutable_data(DDim dims, platform::Place place,
- bool is_pinned) {
+inline T* Tensor::mutable_data(DDim dims, platform::Place place) {
static_assert(std::is_pod::value, "T must be POD");
Resize(dims);
- return mutable_data(place, is_pinned);
+ return mutable_data(place);
}
template
-inline T* Tensor::mutable_data(platform::Place place, bool is_pinned) {
+inline T* Tensor::mutable_data(platform::Place place) {
static_assert(std::is_pod::value, "T must be POD");
- return reinterpret_cast(mutable_data(place, typeid(T), is_pinned));
+ return reinterpret_cast(mutable_data(place, typeid(T)));
}
-inline void* Tensor::mutable_data(platform::Place place, std::type_index type,
- bool is_pinned) {
+inline void* Tensor::mutable_data(platform::Place place, std::type_index type) {
if (holder_ != nullptr) {
holder_->set_type(type);
}
@@ -129,27 +127,26 @@ inline void* Tensor::mutable_data(platform::Place place, std::type_index type,
holder_->size() < size + offset_) {
if (platform::is_cpu_place(place)) {
holder_.reset(new PlaceholderImpl(
- boost::get(place), size, type, is_pinned));
+ boost::get(place), size, type));
} else if (platform::is_gpu_place(place)) {
#ifndef PADDLE_WITH_CUDA
PADDLE_THROW("'CUDAPlace' is not supported in CPU only device.");
}
#else
holder_.reset(new PlaceholderImpl(
- boost::get(place), size, type, is_pinned));
+ boost::get(place), size, type));
}
#endif
offset_ = 0;
- is_pinned_ = is_pinned;
}
return reinterpret_cast(reinterpret_cast(holder_->ptr()) +
offset_);
}
-inline void* Tensor::mutable_data(platform::Place place, bool is_pinned) {
+inline void* Tensor::mutable_data(platform::Place place) {
PADDLE_ENFORCE(this->holder_ != nullptr,
"Cannot invoke mutable data if current hold nothing");
- return mutable_data(place, holder_->type(), is_pinned);
+ return mutable_data(place, holder_->type());
}
inline Tensor& Tensor::ShareDataWith(const Tensor& src) {
@@ -191,8 +188,6 @@ inline const DDim& Tensor::dims() const { return dims_; }
inline int64_t Tensor::numel() const { return product(dims_); }
-inline bool Tensor::isPinned() const { return is_pinned_; }
-
inline Tensor ReshapeToMatrix(const Tensor& src, int num_col_dims) {
Tensor res;
res.ShareDataWith(src);
diff --git a/paddle/fluid/framework/tensor_util.cc b/paddle/fluid/framework/tensor_util.cc
index 8b7533ce71..1d864af011 100644
--- a/paddle/fluid/framework/tensor_util.cc
+++ b/paddle/fluid/framework/tensor_util.cc
@@ -148,6 +148,11 @@ struct AnyVisitor : public boost::static_visitor {
const platform::CPUPlace& cpu) const {
return *out.data();
}
+
+ bool GetResult(const framework::Tensor& out,
+ const platform::CUDAPinnedPlace& cpu) const {
+ return *out.data();
+ }
};
template
diff --git a/paddle/fluid/memory/CMakeLists.txt b/paddle/fluid/memory/CMakeLists.txt
index 1a61c48482..8b3043af7a 100644
--- a/paddle/fluid/memory/CMakeLists.txt
+++ b/paddle/fluid/memory/CMakeLists.txt
@@ -4,13 +4,17 @@ cc_library(memory SRCS memory.cc DEPS place enforce)
cc_library(memcpy SRCS memcpy.cc DEPS place)
cc_library(paddle_memory
- DEPS
- memory
- memcpy
- meta_data
- meta_cache
- memory_block
- buddy_allocator
- system_allocator)
+ DEPS
+ memory
+ memcpy
+ meta_data
+ meta_cache
+ memory_block
+ buddy_allocator
+ system_allocator)
cc_test(memory_test SRCS memory_test.cc DEPS place paddle_memory)
+
+#if (WITH_GPU)
+# nv_test(pinned_memory_test SRCS pinned_memory_test.cu DEPS place paddle_memory)
+#endif()
diff --git a/paddle/fluid/memory/detail/system_allocator.cc b/paddle/fluid/memory/detail/system_allocator.cc
index 22f6f50674..a45f8c33ee 100644
--- a/paddle/fluid/memory/detail/system_allocator.cc
+++ b/paddle/fluid/memory/detail/system_allocator.cc
@@ -14,6 +14,7 @@ limitations under the License. */
#include "paddle/fluid/memory/detail/system_allocator.h"
#include "paddle/fluid/platform/assert.h"
+#include "paddle/fluid/platform/cpu_info.h"
#include "paddle/fluid/platform/enforce.h"
#include "paddle/fluid/platform/gpu_info.h"
@@ -134,21 +135,31 @@ bool GPUAllocator::UseGpu() const { return true; }
// memory. It’s locked to a physical address.
void* CUDAPinnedAllocator::Alloc(size_t& index, size_t size) {
if (size <= 0) return nullptr;
- void* p;
- // NOTE: here, we use GpuMaxAllocSize() as the maximum memory size
+
+ // NOTE: here, we use CUDAPinnedMaxAllocSize as the maximum memory size
// of host pinned allocation. Allocates too much would reduce
// the amount of memory available to the underlying system for paging.
+ size_t usable =
+ paddle::platform::CUDAPinnedMaxAllocSize() - cuda_pinnd_alloc_size_;
- size_t usable = paddle::platform::GpuMaxAllocSize() - fallback_alloc_size_;
-
- if (size > usable) return nullptr;
+ if (size > usable) {
+ LOG(WARNING) << "Cannot malloc " << size / 1024.0 / 1024.0
+ << " MB pinned memory."
+ << ", available " << usable / 1024.0 / 1024.0 << " MB";
+ return nullptr;
+ }
+ void* p;
// PINNED memory is visible to all CUDA contexts.
cudaError_t result = cudaMallocHost(&p, size);
+
if (result == cudaSuccess) {
- index = 1;
- fallback_alloc_size_ += size;
+ index = 1; // PINNED memory
+ cuda_pinnd_alloc_size_ += size;
return p;
+ } else {
+ LOG(WARNING) << "cudaMallocHost failed.";
+ return nullptr;
}
return nullptr;
@@ -158,8 +169,8 @@ void CUDAPinnedAllocator::Free(void* p, size_t size, size_t index) {
cudaError_t err;
PADDLE_ASSERT(index == 1);
- PADDLE_ASSERT(fallback_alloc_size_ >= size);
- fallback_alloc_size_ -= size;
+ PADDLE_ASSERT(cuda_pinnd_alloc_size_ >= size);
+ cuda_pinnd_alloc_size_ -= size;
err = cudaFreeHost(p);
// Purposefully allow cudaErrorCudartUnloading, because
@@ -172,7 +183,7 @@ void CUDAPinnedAllocator::Free(void* p, size_t size, size_t index) {
}
}
-bool CUDAPinnedAllocator::UseGpu() const { return true; }
+bool CUDAPinnedAllocator::UseGpu() const { return false; }
#endif
diff --git a/paddle/fluid/memory/detail/system_allocator.h b/paddle/fluid/memory/detail/system_allocator.h
index e8479e73f4..e3c50ef648 100644
--- a/paddle/fluid/memory/detail/system_allocator.h
+++ b/paddle/fluid/memory/detail/system_allocator.h
@@ -21,8 +21,9 @@ namespace memory {
namespace detail {
/**
- * \brief SystemAllocator is the parent class of CPUAllocator and GPUAllocator.
- * A BuddyAllocator object uses a SystemAllocator* pointing to the
+ * \brief SystemAllocator is the parent class of CPUAllocator,
+ * CUDAPinnedAllocator and GPUAllocator. A BuddyAllocator
+ * object uses a SystemAllocator* pointing to the
* underlying system allocator.
*/
class SystemAllocator {
@@ -62,9 +63,7 @@ class CUDAPinnedAllocator : public SystemAllocator {
virtual bool UseGpu() const;
private:
- size_t gpu_alloc_size_ =
- 0; // TODO(zcd): how to define the upper limit of CUDAPinnedMemory?
- size_t fallback_alloc_size_ = 0;
+ size_t cuda_pinnd_alloc_size_ = 0;
};
#endif
diff --git a/paddle/fluid/memory/memcpy.cc b/paddle/fluid/memory/memcpy.cc
index b991360d04..eddcaab8be 100644
--- a/paddle/fluid/memory/memcpy.cc
+++ b/paddle/fluid/memory/memcpy.cc
@@ -56,6 +56,45 @@ void Copy(
}
}
+template <>
+void Copy(
+ platform::CPUPlace dst_place, void* dst,
+ platform::CUDAPinnedPlace src_place, const void* src, size_t num) {
+ std::memcpy(dst, src, num);
+}
+
+template <>
+void Copy(
+ platform::CUDAPinnedPlace dst_place, void* dst,
+ platform::CPUPlace src_place, const void* src, size_t num) {
+ std::memcpy(dst, src, num);
+}
+
+template <>
+void Copy(
+ platform::CUDAPinnedPlace dst_place, void* dst,
+ platform::CUDAPinnedPlace src_place, const void* src, size_t num) {
+ std::memcpy(dst, src, num);
+}
+
+template <>
+void Copy(
+ platform::CUDAPinnedPlace dst_place, void* dst,
+ platform::CUDAPlace src_place, const void* src, size_t num,
+ cudaStream_t stream) {
+ platform::SetDeviceId(src_place.device);
+ platform::GpuMemcpyAsync(dst, src, num, cudaMemcpyDeviceToHost, stream);
+}
+
+template <>
+void Copy(
+ platform::CUDAPlace dst_place, void* dst,
+ platform::CUDAPinnedPlace src_place, const void* src, size_t num,
+ cudaStream_t stream) {
+ platform::SetDeviceId(dst_place.device);
+ platform::GpuMemcpyAsync(dst, src, num, cudaMemcpyHostToDevice, stream);
+}
+
#endif
} // namespace memory
diff --git a/paddle/fluid/memory/memory.cc b/paddle/fluid/memory/memory.cc
index 56593653a6..09f82166be 100644
--- a/paddle/fluid/memory/memory.cc
+++ b/paddle/fluid/memory/memory.cc
@@ -38,8 +38,7 @@ BuddyAllocator* GetCPUBuddyAllocator() {
}
template <>
-void* Alloc(platform::CPUPlace place, size_t size,
- bool is_pinned) {
+void* Alloc(platform::CPUPlace place, size_t size) {
VLOG(10) << "Allocate " << size << " bytes on " << platform::Place(place);
void* p = GetCPUBuddyAllocator()->Alloc(size);
VLOG(10) << " pointer=" << p;
@@ -47,8 +46,7 @@ void* Alloc(platform::CPUPlace place, size_t size,
}
template <>
-void Free(platform::CPUPlace place, void* p,
- bool is_pinned) {
+void Free(platform::CPUPlace place, void* p) {
VLOG(10) << "Free pointer=" << p << " on " << platform::Place(place);
GetCPUBuddyAllocator()->Free(p);
}
@@ -84,47 +82,15 @@ BuddyAllocator* GetGPUBuddyAllocator(int gpu_id) {
return as[gpu_id];
}
-BuddyAllocator* GetCUDAPinnedBuddyAllocator(int gpu_id) {
- static BuddyAllocator** as = NULL;
- if (as == NULL) {
- int gpu_num = platform::GetCUDADeviceCount();
- as = new BuddyAllocator*[gpu_num];
- for (int gpu = 0; gpu < gpu_num; gpu++) {
- as[gpu] = nullptr;
- }
- }
- platform::SetDeviceId(gpu_id);
- if (!as[gpu_id]) {
- as[gpu_id] = new BuddyAllocator(new detail::CUDAPinnedAllocator,
- platform::GpuMinChunkSize(),
- platform::GpuMaxChunkSize());
- VLOG(10) << "\n\nNOTE: each GPU device use "
- << FLAGS_fraction_of_gpu_memory_to_use * 100
- << "% of GPU memory.\n"
- << "You can set GFlags environment variable '"
- << "FLAGS_fraction_of_gpu_memory_to_use"
- << "' to change the fraction of GPU usage.\n\n";
- }
- return as[gpu_id];
-}
-
template <>
size_t Used(platform::CUDAPlace place) {
return GetGPUBuddyAllocator(place.device)->Used();
}
template <>
-void* Alloc(platform::CUDAPlace place, size_t size,
- bool is_pinned) {
- void* ptr;
- if (is_pinned) {
- auto* buddy_allocator = GetCUDAPinnedBuddyAllocator(place.device);
- ptr = buddy_allocator->Alloc(size);
- } else {
- auto* buddy_allocator = GetGPUBuddyAllocator(place.device);
- ptr = buddy_allocator->Alloc(size);
- }
-
+void* Alloc(platform::CUDAPlace place, size_t size) {
+ auto* buddy_allocator = GetGPUBuddyAllocator(place.device);
+ auto* ptr = buddy_allocator->Alloc(size);
if (ptr == nullptr) {
int cur_dev = platform::GetCurrentDeviceId();
platform::SetDeviceId(place.device);
@@ -142,15 +108,42 @@ void* Alloc(platform::CUDAPlace place, size_t size,
}
template <>
-void Free(platform::CUDAPlace place, void* p,
- bool is_pinned) {
- if (is_pinned) {
- GetCUDAPinnedBuddyAllocator(place.device)->Free(p);
- } else {
- GetGPUBuddyAllocator(place.device)->Free(p);
+void Free(platform::CUDAPlace place, void* p) {
+ GetGPUBuddyAllocator(place.device)->Free(p);
+}
+
+BuddyAllocator* GetCUDAPinnedBuddyAllocator() {
+ static BuddyAllocator* ba = NULL;
+ if (ba == NULL) {
+ ba = new BuddyAllocator(new detail::CUDAPinnedAllocator,
+ platform::CUDAPinnedMinChunkSize(),
+ platform::CUDAPinnedMaxChunkSize());
}
+ return ba;
}
+template <>
+size_t Used(platform::CUDAPinnedPlace place) {
+ return GetCUDAPinnedBuddyAllocator()->Used();
+}
+
+template <>
+void* Alloc(platform::CUDAPinnedPlace place,
+ size_t size) {
+ auto* buddy_allocator = GetCUDAPinnedBuddyAllocator();
+ void* ptr = buddy_allocator->Alloc(size);
+
+ if (ptr == nullptr) {
+ LOG(WARNING) << "cudaMallocHost Cannot allocate " << size
+ << " bytes in CUDAPinnedPlace";
+ }
+ return ptr;
+}
+
+template <>
+void Free(platform::CUDAPinnedPlace place, void* p) {
+ GetCUDAPinnedBuddyAllocator()->Free(p);
+}
#endif
size_t Usage::operator()(const platform::CPUPlace& cpu) const {
@@ -165,6 +158,14 @@ size_t Usage::operator()(const platform::CUDAPlace& gpu) const {
#endif
}
+size_t Usage::operator()(const platform::CUDAPinnedPlace& cuda_pinned) const {
+#ifdef PADDLE_WITH_CUDA
+ return Used(cuda_pinned);
+#else
+ PADDLE_THROW("'CUDAPinnedPlace' is not supported in CPU only device.");
+#endif
+}
+
size_t memory_usage(const platform::Place& p) {
return boost::apply_visitor(Usage(), p);
}
diff --git a/paddle/fluid/memory/memory.h b/paddle/fluid/memory/memory.h
index 062bfc880e..3e6bfddd69 100644
--- a/paddle/fluid/memory/memory.h
+++ b/paddle/fluid/memory/memory.h
@@ -33,7 +33,7 @@ namespace memory {
* address is valid or not.
*/
template
-void* Alloc(Place place, size_t size, bool is_pinned = false);
+void* Alloc(Place place, size_t size);
/**
* \brief Free memory block in one place.
@@ -43,7 +43,7 @@ void* Alloc(Place place, size_t size, bool is_pinned = false);
*
*/
template
-void Free(Place place, void* ptr, bool is_pinned = false);
+void Free(Place place, void* ptr);
/**
* \brief Total size of used memory in one place.
@@ -57,6 +57,7 @@ size_t Used(Place place);
struct Usage : public boost::static_visitor {
size_t operator()(const platform::CPUPlace& cpu) const;
size_t operator()(const platform::CUDAPlace& gpu) const;
+ size_t operator()(const platform::CUDAPinnedPlace& cuda_pinned) const;
};
size_t memory_usage(const platform::Place& p);
@@ -74,13 +75,11 @@ class PODDeleter {
static_assert(std::is_pod::value, "T must be POD");
public:
- explicit PODDeleter(Place place, bool is_pinned = false)
- : place_(place), is_pinned_(is_pinned) {}
- void operator()(T* ptr) { Free(place_, static_cast(ptr), is_pinned_); }
+ explicit PODDeleter(Place place) : place_(place) {}
+ void operator()(T* ptr) { Free(place_, static_cast(ptr)); }
private:
Place place_;
- bool is_pinned_;
};
/**
diff --git a/paddle/fluid/memory/memory_test.cc b/paddle/fluid/memory/memory_test.cc
index eb27a52b25..03829702a0 100644
--- a/paddle/fluid/memory/memory_test.cc
+++ b/paddle/fluid/memory/memory_test.cc
@@ -141,4 +141,59 @@ TEST(BuddyAllocator, GPUMultAlloc) {
}
}
+size_t align(size_t size, paddle::platform::CUDAPinnedPlace place) {
+ size += sizeof(paddle::memory::detail::Metadata);
+ size_t alignment = paddle::platform::CUDAPinnedMinChunkSize();
+ size_t remaining = size % alignment;
+ return remaining == 0 ? size : size + (alignment - remaining);
+}
+
+TEST(BuddyAllocator, CUDAPinnedAllocator) {
+ void *p = nullptr;
+
+ EXPECT_EQ(p, nullptr);
+
+ paddle::platform::CUDAPinnedPlace cpu;
+ p = paddle::memory::Alloc(cpu, 4096);
+
+ EXPECT_NE(p, nullptr);
+
+ paddle::platform::Place place = cpu;
+ EXPECT_EQ(paddle::memory::Used(cpu), paddle::memory::memory_usage(place));
+
+ paddle::memory::Free(cpu, p);
+}
+
+TEST(BuddyAllocator, CUDAPinnedMultAllocator) {
+ paddle::platform::CUDAPinnedPlace cpu;
+
+ std::unordered_map ps;
+
+ size_t total_size = paddle::memory::Used(cpu);
+ EXPECT_EQ(total_size, 0UL);
+
+ for (auto size :
+ {0, 128, 256, 1024, 4096, 16384, 65536, 262144, 1048576, 4194304}) {
+ ps[paddle::memory::Alloc(cpu, size)] = size;
+
+ // Buddy Allocator doesn't manage too large memory chunk
+ if (paddle::memory::Used(cpu) == total_size) continue;
+
+ size_t aligned_size = align(size, cpu);
+ total_size += aligned_size;
+ EXPECT_EQ(total_size, paddle::memory::Used(cpu));
+ }
+
+ for (auto p : ps) {
+ EXPECT_EQ(is_aligned(p.first), true);
+ paddle::memory::Free(cpu, p.first);
+
+ // Buddy Allocator doesn't manage too large memory chunk
+ if (paddle::memory::Used(cpu) == total_size) continue;
+
+ size_t aligned_size = align(p.second, cpu);
+ total_size -= aligned_size;
+ EXPECT_EQ(total_size, paddle::memory::Used(cpu));
+ }
+}
#endif
diff --git a/paddle/fluid/memory/pinned_memory_test.cu b/paddle/fluid/memory/pinned_memory_test.cu
new file mode 100644
index 0000000000..a000001f41
--- /dev/null
+++ b/paddle/fluid/memory/pinned_memory_test.cu
@@ -0,0 +1,147 @@
+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+
+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
+#include
+
+#include "paddle/fluid/memory/detail/memory_block.h"
+#include "paddle/fluid/memory/detail/meta_data.h"
+#include "paddle/fluid/memory/memcpy.h"
+#include "paddle/fluid/memory/memory.h"
+
+#include "paddle/fluid/platform/cpu_info.h"
+#include "paddle/fluid/platform/gpu_info.h"
+#include "paddle/fluid/platform/place.h"
+
+// This unit test is an example comparing the performance between using pinned
+// memory and not. In general, using pinned memory will be faster.
+template
+__global__ void Kernel(T* output, int dim) {
+ int tid = blockIdx.x * blockDim.x + threadIdx.x;
+ if (tid < dim) {
+ output[tid] = output[tid] * output[tid] / 100;
+ }
+}
+
+template
+float test_pinned_memory() {
+ Place cpu_place;
+ paddle::platform::CUDAPlace cuda_place;
+
+ const int data_size = 4096;
+ const int iteration = 10;
+
+ // create event start and end
+ cudaEvent_t start_e, stop_e, copying_e;
+ float elapsedTime = 0;
+ cudaEventCreate(&start_e);
+ cudaEventCreate(&stop_e);
+ cudaEventCreate(©ing_e);
+
+ // create computation stream, data copying stream
+ cudaStream_t computation_stream, copying_stream;
+ cudaStreamCreate(&computation_stream);
+ cudaStreamCreate(©ing_stream);
+
+ // create record event, pinned memory, gpu memory
+ std::vector record_event(iteration);
+ std::vector input_pinned_mem(iteration);
+ std::vector gpu_mem(iteration);
+ std::vector output_pinned_mem(iteration);
+
+ // initial data
+ for (int j = 0; j < iteration; ++j) {
+ cudaEventCreateWithFlags(&record_event[j], cudaEventDisableTiming);
+ cudaEventCreate(&(record_event[j]));
+ input_pinned_mem[j] = static_cast(
+ paddle::memory::Alloc(cpu_place, data_size * sizeof(float)));
+ output_pinned_mem[j] = static_cast(
+ paddle::memory::Alloc(cpu_place, data_size * sizeof(float)));
+ gpu_mem[j] = static_cast(
+ paddle::memory::Alloc(cuda_place, data_size * sizeof(float)));
+
+ for (int k = 0; k < data_size; ++k) {
+ input_pinned_mem[j][k] = k;
+ }
+ }
+
+ cudaEventRecord(start_e, computation_stream);
+
+ // computation
+ for (int m = 0; m < 30; ++m) {
+ for (int i = 0; i < iteration; ++i) {
+ // cpu -> GPU on computation stream.
+ // note: this operation is async for pinned memory.
+ paddle::memory::Copy(cuda_place, gpu_mem[i], cpu_place,
+ input_pinned_mem[i], data_size * sizeof(float),
+ computation_stream);
+
+ // call kernel on computation stream.
+ Kernel<<<4, 1024, 0, computation_stream>>>(gpu_mem[i], data_size);
+
+ // record event_computation on computation stream
+ cudaEventRecord(record_event[i], computation_stream);
+
+ // wait event_computation on copy stream.
+ // note: this operation is async.
+ cudaStreamWaitEvent(copying_stream, record_event[i], 0);
+
+ // copy data GPU->CPU, on copy stream.
+ // note: this operation is async for pinned memory.
+ paddle::memory::Copy(cpu_place, output_pinned_mem[i], cuda_place,
+ gpu_mem[i], data_size * sizeof(float),
+ copying_stream);
+ }
+ }
+
+ cudaEventRecord(copying_e, copying_stream);
+ cudaStreamWaitEvent(computation_stream, copying_e, 0);
+
+ cudaEventRecord(stop_e, computation_stream);
+
+ cudaEventSynchronize(start_e);
+ cudaEventSynchronize(stop_e);
+ cudaEventElapsedTime(&elapsedTime, start_e, stop_e);
+
+ // std::cout << cpu_place << " "
+ // << "time consume:" << elapsedTime / 30 << std::endl;
+
+ for (int l = 0; l < iteration; ++l) {
+ for (int k = 0; k < data_size; ++k) {
+ float temp = input_pinned_mem[l][k];
+ temp = temp * temp / 100;
+ EXPECT_FLOAT_EQ(temp, output_pinned_mem[l][k]);
+ }
+ }
+
+ // destroy resource
+ cudaEventDestroy(copying_e);
+ cudaEventDestroy(start_e);
+ cudaEventDestroy(stop_e);
+ for (int j = 0; j < 10; ++j) {
+ cudaEventDestroy((record_event[j]));
+ paddle::memory::Free(cpu_place, input_pinned_mem[j]);
+ paddle::memory::Free(cpu_place, output_pinned_mem[j]);
+ paddle::memory::Free(cuda_place, gpu_mem[j]);
+ }
+ return elapsedTime / 30;
+}
+
+TEST(CPUANDCUDAPinned, CPUAllocatorAndCUDAPinnedAllocator) {
+ // Generally speaking, operation on pinned_memory is faster than that on
+ // unpinned-memory, but if this unit test fails frequently, please close this
+ // test for the time being.
+ float time1 = test_pinned_memory();
+ float time2 = test_pinned_memory();
+ EXPECT_GT(time1, time2);
+}
diff --git a/paddle/fluid/operators/CMakeLists.txt b/paddle/fluid/operators/CMakeLists.txt
index 9ed79453b9..952ac8b1dc 100644
--- a/paddle/fluid/operators/CMakeLists.txt
+++ b/paddle/fluid/operators/CMakeLists.txt
@@ -193,6 +193,7 @@ if(WITH_DISTRIBUTE)
set_source_files_properties(send_vars_op.cc PROPERTIES COMPILE_FLAGS ${DISTRIBUTE_COMPILE_FLAGS})
op_library(send_barrier_op DEPS ${DISTRIBUTE_DEPS})
set_source_files_properties(send_barrier_op.cc PROPERTIES COMPILE_FLAGS ${DISTRIBUTE_COMPILE_FLAGS})
+ set_source_files_properties(send_recv_op_test.cc PROPERTIES COMPILE_FLAGS ${DISTRIBUTE_COMPILE_FLAGS})
cc_test(test_send_recv SRCS send_recv_op_test.cc DEPS prefetch_op send_op listen_and_serv_op sum_op executor)
else()
set(DEPS_OPS ${DEPS_OPS} send_op prefetch_op recv_op listen_and_serv_op send_vars_op send_barrier_op)
diff --git a/paddle/fluid/operators/conv_op.cc b/paddle/fluid/operators/conv_op.cc
index 650bc92be2..695db841a4 100644
--- a/paddle/fluid/operators/conv_op.cc
+++ b/paddle/fluid/operators/conv_op.cc
@@ -13,6 +13,10 @@ See the License for the specific language governing permissions and
limitations under the License. */
#include "paddle/fluid/operators/conv_op.h"
+
+#include
+#include
+
#ifdef PADDLE_WITH_CUDA
#include "paddle/fluid/platform/cudnn_helper.h"
#endif
diff --git a/paddle/fluid/operators/detail/bytebuffer_stream.cc b/paddle/fluid/operators/detail/bytebuffer_stream.cc
index 741dd51de9..a14171563e 100644
--- a/paddle/fluid/operators/detail/bytebuffer_stream.cc
+++ b/paddle/fluid/operators/detail/bytebuffer_stream.cc
@@ -17,7 +17,7 @@ limitations under the License. */
// file and did some modifications so that we can send gRPC
// requests without too much copying of the tensor data.
-#include "bytebuffer_stream.h"
+#include "paddle/fluid/operators/detail/bytebuffer_stream.h"
namespace paddle {
namespace operators {
diff --git a/paddle/fluid/operators/detail/bytebuffer_stream.h b/paddle/fluid/operators/detail/bytebuffer_stream.h
index 1791a48aab..054dd4ff29 100644
--- a/paddle/fluid/operators/detail/bytebuffer_stream.h
+++ b/paddle/fluid/operators/detail/bytebuffer_stream.h
@@ -19,9 +19,11 @@ limitations under the License. */
#pragma once
-#include
+#include
+
#include "google/protobuf/io/coded_stream.h"
#include "google/protobuf/io/zero_copy_stream.h"
+#include "grpc++/grpc++.h"
namespace grpc {
// A ZeroCopyInputStream that reads from grpc_byte_buffer
@@ -56,7 +58,7 @@ class GrpcBufferReader final
*data = GRPC_SLICE_START_PTR(slice_) + GRPC_SLICE_LENGTH(slice_) -
backup_count_;
GPR_CODEGEN_ASSERT(backup_count_ <= INT_MAX);
- *size = (int)backup_count_;
+ *size = static_cast(backup_count_);
backup_count_ = 0;
return true;
}
@@ -68,7 +70,7 @@ class GrpcBufferReader final
*data = GRPC_SLICE_START_PTR(slice_);
// On win x64, int is only 32bit
GPR_CODEGEN_ASSERT(GRPC_SLICE_LENGTH(slice_) <= INT_MAX);
- byte_count_ += * size = (int)GRPC_SLICE_LENGTH(slice_);
+ byte_count_ += * size = static_cast(GRPC_SLICE_LENGTH(slice_));
return true;
}
diff --git a/paddle/fluid/operators/detail/grpc_client.cc b/paddle/fluid/operators/detail/grpc_client.cc
index d79ba6d291..ef987d07f0 100644
--- a/paddle/fluid/operators/detail/grpc_client.cc
+++ b/paddle/fluid/operators/detail/grpc_client.cc
@@ -14,6 +14,8 @@ limitations under the License. */
#include "paddle/fluid/operators/detail/grpc_client.h"
+#include
+
#include
#include "paddle/fluid/framework/threadpool.h"
@@ -54,7 +56,7 @@ bool RPCClient::AsyncSendVariable(const std::string& ep,
auto call = s->stub_g_.PrepareUnaryCall(
s->context_.get(), "/sendrecv.SendRecvService/SendVariable", req, &cq_);
call->StartCall();
- call->Finish(&s->reply_, &s->status_, static_cast(s));
+ call->Finish(&s->reply_, &s->status_, reinterpret_cast(s));
});
req_count_++;
@@ -66,7 +68,7 @@ void ProcGetResponse(const VarHandle& var_h,
// const sendrecv::VariableMessage& ret_msg) {
const ::grpc::ByteBuffer& ret_msg) {
framework::Variable* outvar = NULL;
- DeserializeFromByteBuffer(ret_msg, *var_h.ctx, var_h.scope, outvar);
+ DeserializeFromByteBuffer(ret_msg, *var_h.ctx, var_h.scope, &outvar);
}
template
@@ -110,7 +112,7 @@ bool RPCClient::AsyncGetVariable(const std::string& ep,
auto call = s->stub_g_.PrepareUnaryCall(
s->context_.get(), "/sendrecv.SendRecvService/GetVariable", buf, &cq_);
call->StartCall();
- call->Finish(&s->reply_, &s->status_, static_cast(s));
+ call->Finish(&s->reply_, &s->status_, reinterpret_cast(s));
});
req_count_++;
@@ -170,7 +172,7 @@ void RPCClient::AsyncSendBatchBarrier(const std::string& ep, int64_t time_out) {
sendrecv::VariableMessage req;
req.set_varname(BATCH_BARRIER_MESSAGE);
auto rpc = s->stub_->AsyncSendVariable(s->context_.get(), req, &cq_);
- rpc->Finish(&s->reply_, &s->status_, static_cast(s));
+ rpc->Finish(&s->reply_, &s->status_, reinterpret_cast(s));
req_count_++;
}
@@ -182,7 +184,7 @@ void RPCClient::AsyncSendFetchBarrier(const std::string& ep, int64_t time_out) {
sendrecv::VariableMessage req;
req.set_varname(FETCH_BARRIER_MESSAGE);
auto rpc = s->stub_->AsyncGetVariable(s->context_.get(), req, &cq_);
- rpc->Finish(&s->reply_, &s->status_, static_cast(s));
+ rpc->Finish(&s->reply_, &s->status_, reinterpret_cast(s));
req_count_++;
}
diff --git a/paddle/fluid/operators/detail/grpc_client.h b/paddle/fluid/operators/detail/grpc_client.h
index fe237e54ef..4425b19328 100644
--- a/paddle/fluid/operators/detail/grpc_client.h
+++ b/paddle/fluid/operators/detail/grpc_client.h
@@ -14,10 +14,9 @@ limitations under the License. */
#pragma once
-#include
-#include
#include
-#include
+
+#include // NOLINT
#include
#include
#include
@@ -25,11 +24,11 @@ limitations under the License. */
#include
#include
-#include
-#include
-#include
-#include
-
+#include "grpc++/generic/generic_stub.h"
+#include "grpc++/grpc++.h"
+#include "grpc++/support/byte_buffer.h"
+#include "grpc++/support/slice.h"
+#include "grpc/support/log.h"
#include "paddle/fluid/framework/data_type.h"
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/scope.h"
diff --git a/paddle/fluid/operators/detail/grpc_server.cc b/paddle/fluid/operators/detail/grpc_server.cc
index 7c978b28b6..2e7bf1921a 100644
--- a/paddle/fluid/operators/detail/grpc_server.cc
+++ b/paddle/fluid/operators/detail/grpc_server.cc
@@ -186,7 +186,8 @@ void AsyncGRPCServer::WaitClientGet(int count) {
void AsyncGRPCServer::RunSyncUpdate() {
::grpc::ServerBuilder builder;
- builder.AddListeningPort(address_, ::grpc::InsecureServerCredentials());
+ builder.AddListeningPort(address_, ::grpc::InsecureServerCredentials(),
+ &selected_port_);
builder.SetMaxSendMessageSize(std::numeric_limits::max());
builder.SetMaxReceiveMessageSize(std::numeric_limits::max());
builder.RegisterService(&service_);
@@ -196,7 +197,8 @@ void AsyncGRPCServer::RunSyncUpdate() {
cq_prefetch_ = builder.AddCompletionQueue();
server_ = builder.BuildAndStart();
- LOG(INFO) << "Server listening on " << address_ << std::endl;
+ LOG(INFO) << "Server listening on " << address_
+ << " selected port: " << selected_port_;
std::function send_register =
std::bind(&AsyncGRPCServer::TryToRegisterNewSendOne, this);
@@ -273,7 +275,7 @@ void AsyncGRPCServer::TryToRegisterNewPrefetchOne() {
// FIXME(typhoonzero): change cq_name to enum.
void AsyncGRPCServer::HandleRequest(::grpc::ServerCompletionQueue* cq,
- std::string cq_name,
+ const std::string& cq_name,
std::function TryToRegisterNewOne) {
TryToRegisterNewOne();
diff --git a/paddle/fluid/operators/detail/grpc_server.h b/paddle/fluid/operators/detail/grpc_server.h
index b0596d3cd1..380447f47c 100644
--- a/paddle/fluid/operators/detail/grpc_server.h
+++ b/paddle/fluid/operators/detail/grpc_server.h
@@ -14,10 +14,11 @@ limitations under the License. */
#pragma once
-#include
#include
+#include // NOLINT
#include
+#include "grpc++/grpc++.h"
#include "paddle/fluid/framework/executor.h"
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/program_desc.h"
@@ -62,6 +63,8 @@ class AsyncGRPCServer final {
void SetExecutor(framework::Executor *executor) { executor_ = executor; }
+ int GetSelectedPort() { return selected_port_; }
+
const ReceivedMessage Get() { return this->var_recv_queue_.Pop(); }
void Push(const std::string &msg_name) {
@@ -71,7 +74,8 @@ class AsyncGRPCServer final {
void ShutDown();
protected:
- void HandleRequest(::grpc::ServerCompletionQueue *cq, std::string cq_name,
+ void HandleRequest(::grpc::ServerCompletionQueue *cq,
+ const std::string &cq_name,
std::function TryToRegisterNewOne);
void TryToRegisterNewSendOne();
void TryToRegisterNewGetOne();
@@ -109,6 +113,7 @@ class AsyncGRPCServer final {
int prefetch_blk_id_;
framework::ProgramDesc *program_;
framework::Executor *executor_;
+ int selected_port_;
};
}; // namespace detail
diff --git a/paddle/fluid/operators/detail/grpc_server_test.cc b/paddle/fluid/operators/detail/grpc_server_test.cc
index 1ad62863a1..b89aed0157 100644
--- a/paddle/fluid/operators/detail/grpc_server_test.cc
+++ b/paddle/fluid/operators/detail/grpc_server_test.cc
@@ -14,7 +14,7 @@ limitations under the License. */
#include
#include
-#include
+#include // NOLINT
#include "gtest/gtest.h"
#include "paddle/fluid/operators/detail/grpc_client.h"
diff --git a/paddle/fluid/operators/detail/proto_encoder_helper.h b/paddle/fluid/operators/detail/proto_encoder_helper.h
index 4a7bfb8bd5..d91d054b25 100644
--- a/paddle/fluid/operators/detail/proto_encoder_helper.h
+++ b/paddle/fluid/operators/detail/proto_encoder_helper.h
@@ -19,7 +19,9 @@ limitations under the License. */
#pragma once
-#include
+#include
+
+#include "grpc++/grpc++.h"
#include "paddle/fluid/platform/enforce.h"
namespace paddle {
@@ -142,6 +144,6 @@ class ProtoEncodeHelper {
char* limit_; // Just for CHECKs
};
-} // detail
-} // operators
-} // paddle
+} // namespace detail
+} // namespace operators
+} // namespace paddle
diff --git a/paddle/fluid/operators/detail/sendrecvop_utils.cc b/paddle/fluid/operators/detail/sendrecvop_utils.cc
index 7e3f015dab..f8576d01b1 100644
--- a/paddle/fluid/operators/detail/sendrecvop_utils.cc
+++ b/paddle/fluid/operators/detail/sendrecvop_utils.cc
@@ -13,8 +13,10 @@ See the License for the specific language governing permissions and
limitations under the License. */
#include "paddle/fluid/operators/detail/sendrecvop_utils.h"
+
#include
-#include
+#include // NOLINT
+
#include "google/protobuf/io/coded_stream.h"
#include "google/protobuf/io/zero_copy_stream.h"
#include "paddle/fluid/framework/data_type.h"
@@ -42,7 +44,7 @@ void SerializeToByteBuffer(const std::string& name, framework::Variable* var,
void* buf = malloc(1024);
void* payload = nullptr;
size_t payload_size;
- ProtoEncodeHelper e((char*)buf, 1024);
+ ProtoEncodeHelper e(static_cast(buf), 1024);
e.WriteString(VarMsg::kVarnameFieldNumber, name);
if (var->IsType()) {
e.WriteUint64(VarMsg::kTypeFieldNumber, 0);
@@ -152,7 +154,7 @@ void SerializeToByteBuffer(const std::string& name, framework::Variable* var,
framework::proto::VarType_Type_SELECTED_ROWS) {
auto* slr = var->GetMutable();
- ProtoEncodeHelper e2((char*)buf, 128);
+ ProtoEncodeHelper e2(static_cast(buf), 128);
// NOTE: rows is of type int64_t
size_t rows_memory_size =
slr->rows().size() * framework::SizeOfType(typeid(int64_t));
@@ -181,10 +183,10 @@ void SerializeToByteBuffer(const std::string& name, framework::Variable* var,
void DeserializeFromByteBuffer(const ::grpc::ByteBuffer& msg,
const platform::DeviceContext& ctx,
const framework::Scope* scope,
- framework::Variable*& var) {
+ framework::Variable** var) {
operators::detail::VariableResponse resp(scope, &ctx);
PADDLE_ENFORCE(resp.Parse(msg) == 0, "parse bytebuffer to tensor error!");
- var = resp.GetVar();
+ *var = resp.GetVar();
}
} // namespace detail
diff --git a/paddle/fluid/operators/detail/sendrecvop_utils.h b/paddle/fluid/operators/detail/sendrecvop_utils.h
index b3b2b8469c..d795444084 100644
--- a/paddle/fluid/operators/detail/sendrecvop_utils.h
+++ b/paddle/fluid/operators/detail/sendrecvop_utils.h
@@ -51,7 +51,7 @@ void SerializeToByteBuffer(const std::string& name, framework::Variable* var,
void DeserializeFromByteBuffer(const ::grpc::ByteBuffer& msg,
const platform::DeviceContext& ctx,
const framework::Scope* scope,
- framework::Variable*& var);
+ framework::Variable** var);
inline std::type_index ToTypeIndex(sendrecv::VariableMessage::Type type) {
switch (type) {
diff --git a/paddle/fluid/operators/detail/serde_test.cc b/paddle/fluid/operators/detail/serde_test.cc
index ea1670e56f..f8cae6b26a 100644
--- a/paddle/fluid/operators/detail/serde_test.cc
+++ b/paddle/fluid/operators/detail/serde_test.cc
@@ -14,9 +14,9 @@ limitations under the License. */
#include
#include
-#include
+#include // NOLINT
-#include
+#include "google/protobuf/text_format.h"
#include "gtest/gtest.h"
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/tensor_util.h"
@@ -107,7 +107,7 @@ void RunSerdeTestSelectedRows(platform::Place place) {
for (int i = 0; i < tensor_numel; ++i) {
EXPECT_FLOAT_EQ(tensor_data2[i], 32.7);
}
- for (int i = 0; i < rows2->size(); ++i) {
+ for (int64_t i = 0; i < rows2->size(); ++i) {
EXPECT_EQ(rows_data2[i], i);
}
EXPECT_EQ(slr2->height(), 1000);
diff --git a/paddle/fluid/operators/detail/simple_block_queue.h b/paddle/fluid/operators/detail/simple_block_queue.h
index 36b58b0c67..69773e05df 100644
--- a/paddle/fluid/operators/detail/simple_block_queue.h
+++ b/paddle/fluid/operators/detail/simple_block_queue.h
@@ -14,9 +14,9 @@ limitations under the License. */
#pragma once
-#include
+#include // NOLINT
#include
-#include
+#include // NOLINT
namespace paddle {
namespace operators {
diff --git a/paddle/fluid/operators/detail/variable_response.cc b/paddle/fluid/operators/detail/variable_response.cc
index f59c9b50bb..78e1d274a9 100644
--- a/paddle/fluid/operators/detail/variable_response.cc
+++ b/paddle/fluid/operators/detail/variable_response.cc
@@ -13,7 +13,11 @@
// limitations under the License.
#include "paddle/fluid/operators/detail/variable_response.h"
-#include
+
+#include
+#include
+#include
+
#include "paddle/fluid/operators/detail/send_recv.pb.h"
#include "paddle/fluid/operators/detail/sendrecvop_utils.h"
@@ -108,7 +112,8 @@ bool ReadRaw(::google::protobuf::io::CodedInputStream* input,
bool VariableResponse::CopyLodTensorData(
::google::protobuf::io::CodedInputStream* input,
- const platform::DeviceContext& ctx, framework::DDim& dims, int length) {
+ const platform::DeviceContext& ctx, const framework::DDim& dims,
+ int length) {
auto var = scope_->FindVar(meta_.varname());
auto* tensor = var->GetMutable();
tensor->Resize(dims);
@@ -144,14 +149,15 @@ inline framework::DDim GetDims(
bool VariableResponse::CopySelectRowsTensorData(
::google::protobuf::io::CodedInputStream* input,
- const platform::DeviceContext& ctx, framework::DDim& dims, int length) {
+ const platform::DeviceContext& ctx, const framework::DDim& dims,
+ int length) {
auto var = scope_->FindVar(meta_.varname());
auto* slr = var->GetMutable();
slr->set_height(meta_.slr_height());
auto* tensor = slr->mutable_value();
tensor->Resize(dims);
PADDLE_ENFORCE_EQ(
- tensor->numel(),
+ static_cast(tensor->numel()),
length / framework::SizeOfType(
paddle::operators::detail::ToTypeIndex(meta_.data_type())));
void* tensor_data = tensor->mutable_data(
diff --git a/paddle/fluid/operators/detail/variable_response.h b/paddle/fluid/operators/detail/variable_response.h
index e121ed7bce..050b6b8401 100644
--- a/paddle/fluid/operators/detail/variable_response.h
+++ b/paddle/fluid/operators/detail/variable_response.h
@@ -14,6 +14,8 @@
#pragma once
+#include
+
#include "paddle/fluid/framework/data_type.h"
#include "paddle/fluid/framework/lod_tensor.h"
#include "paddle/fluid/framework/scope.h"
@@ -60,14 +62,14 @@ class VariableResponse {
private:
bool CopySelectRowsTensorData(::google::protobuf::io::CodedInputStream* input,
const platform::DeviceContext& ctx,
- framework::DDim& dims, int length);
+ const framework::DDim& dims, int length);
bool CopySelectRowsData(::google::protobuf::io::CodedInputStream* input,
const platform::DeviceContext& ctx, int length);
bool CopyLodTensorData(::google::protobuf::io::CodedInputStream* input,
const platform::DeviceContext& ctx,
- framework::DDim& dims, int length);
+ const framework::DDim& dims, int length);
private:
const framework::Scope* scope_;
diff --git a/paddle/fluid/operators/fc_mkldnn_op.cc b/paddle/fluid/operators/fc_mkldnn_op.cc
new file mode 100644
index 0000000000..9c704a2949
--- /dev/null
+++ b/paddle/fluid/operators/fc_mkldnn_op.cc
@@ -0,0 +1,303 @@
+/* Copyright (c) 2018 PaddlePaddle Authors. All Rights Reserved.
+
+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/fluid/framework/tensor.h"
+#include "paddle/fluid/operators/fc_op.h"
+#include "paddle/fluid/platform/device_context.h"
+#include "paddle/fluid/platform/mkldnn_helper.h"
+
+namespace paddle {
+namespace operators {
+
+using paddle::framework::Tensor;
+using paddle::platform::MKLDNNDeviceContext;
+
+template
+class MKLDNNMD {
+ public:
+ explicit MKLDNNMD(const T* in, const T* w, bool bias)
+ : in{paddle::framework::vectorize2int(in->dims())},
+ w{paddle::framework::vectorize2int(w->dims())} {
+ with_bias_ = bias;
+ }
+
+ mkldnn::memory::desc dst() const {
+ return platform::MKLDNNMemDesc({in[0], w[1]},
+ mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::nc);
+ }
+
+ mkldnn::memory::desc src() const {
+ return is_spatial()
+ ? platform::MKLDNNMemDesc({in[0], in[1], in[2], in[3]},
+ mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::nchw)
+ : platform::MKLDNNMemDesc({in[0], in[1]},
+ mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::nc);
+ }
+
+ mkldnn::memory::desc weights() const {
+ return is_spatial()
+ ? platform::MKLDNNMemDesc({w[1], in[1], in[2], in[3]},
+ mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::oihw)
+ : platform::MKLDNNMemDesc({w[1], in[1]},
+ mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::oi);
+ }
+
+ mkldnn::memory::desc bias() const {
+ return with_bias_
+ ? platform::MKLDNNMemDesc({w[1]}, mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::format_undef)
+ : platform::MKLDNNMemDesc({}, mkldnn::memory::data_type::f32,
+ mkldnn::memory::format::format_undef);
+ }
+
+ private:
+ bool is_spatial() const { return in.size() > 1 && w.size() > 1; }
+
+ std::vector in;
+ std::vector w;
+ bool with_bias_;
+ bool is_spatial_;
+};
+
+class MKLDNNMemory {
+ public:
+ MKLDNNMemory(MKLDNNMD* t, const mkldnn::engine& e)
+ : md_{t}, engine_{e} {}
+ virtual ~MKLDNNMemory() = default;
+
+ template
+ mkldnn::memory dst(const Output* out) {
+ return mkldnn::memory({md_->dst(), engine_},
+ static_cast(const_cast(out)));
+ }
+
+ template
+ mkldnn::memory dst(Output* out) {
+ return mkldnn::memory({md_->dst(), engine_}, out);
+ }
+
+ template
+ mkldnn::memory src(const Input* in) {
+ return mkldnn::memory({md_->src(), engine_},
+ static_cast(const_cast(in)));
+ }
+
+ template
+ mkldnn::memory weights(const Weight* w) {
+ return mkldnn::memory({md_->weights(), engine_},
+ static_cast(const_cast(w)));
+ }
+
+ mkldnn::memory bias() {
+ return mkldnn::memory(mkldnn::memory::primitive_desc(md_->bias(), engine_));
+ }
+
+ private:
+ MKLDNNMD* md_;
+ const mkldnn::engine& engine_;
+};
+
+template
+class FCMKLDNNOpKernel : public paddle::framework::OpKernel {
+ void Compute(const paddle::framework::ExecutionContext& ctx) const override {
+ PADDLE_ENFORCE(paddle::platform::is_cpu_place(ctx.GetPlace()),
+ "It must use CPUPlace.");
+
+ auto& dev_ctx = ctx.template device_context();
+ const auto& mkldnn_engine = dev_ctx.GetEngine();
+
+ auto input = ctx.Input("Input");
+ auto w = ctx.Input("W");
+
+ PADDLE_ENFORCE(input->dims().size() == 2 || input->dims().size() == 4,
+ "Input must be with 2 or 4 dimensions, i.e. NCHW");
+ PADDLE_ENFORCE(w->dims().size() == 2 || w->dims().size() == 4,
+ "Weights must be with 2 or 4 dimensions, i.e. OI or OIHW");
+
+ bool with_bias = ctx.Attr("bias_attr");
+ MKLDNNMD md(input, w, with_bias);
+
+ std::shared_ptr pd =
+ FcFwdPrimitiveDesc(md.src(), md.weights(), md.dst(), md.bias(),
+ with_bias, mkldnn_engine);
+
+ const std::string key = ctx.op().Output("Out");
+ const std::string key_fc_pd = key + "@fc_pd";
+
+ dev_ctx.SetBlob(key_fc_pd, pd);
+
+ MKLDNNMemory mem(&md, mkldnn_engine);
+
+ const T* input_data = input->data();
+ const T* w_data = w->data();
+
+ auto output = ctx.Output("Out");
+ T* output_data = output->mutable_data(ctx.GetPlace());
+
+ auto dst_memory = mem.dst(output_data);
+ auto src_memory = mem.src(input_data);
+ auto weights_memory = mem.weights(w_data);
+ auto bias_memory = mem.bias();
+
+ auto forward = with_bias ? mkldnn::inner_product_forward(
+ *pd, src_memory, weights_memory, bias_memory,
+ dst_memory)
+ : mkldnn::inner_product_forward(
+ *pd, src_memory, weights_memory, dst_memory);
+
+ std::vector pipeline = {forward};
+ mkldnn::stream(mkldnn::stream::kind::eager).submit(pipeline).wait();
+ }
+
+ private:
+ std::unique_ptr
+ FcFwdPrimitiveDesc(const mkldnn::memory::desc& src,
+ const mkldnn::memory::desc& weights,
+ const mkldnn::memory::desc& dst,
+ const mkldnn::memory::desc& bias, const bool with_bias,
+ const mkldnn::engine& engine) const {
+ auto desc = with_bias
+ ? mkldnn::inner_product_forward::desc(
+ mkldnn::prop_kind::forward, src, weights, bias, dst)
+ : mkldnn::inner_product_forward::desc(
+ mkldnn::prop_kind::forward, src, weights, dst);
+
+ auto pd = new mkldnn::inner_product_forward::primitive_desc(desc, engine);
+ return std::unique_ptr(pd);
+ }
+};
+
+template
+class FCMKLDNNGradOpKernel : public paddle::framework::OpKernel {
+ public:
+ void Compute(const paddle::framework::ExecutionContext& ctx) const override {
+ PADDLE_ENFORCE(paddle::platform::is_cpu_place(ctx.GetPlace()),
+ "It must use CPUPlace.");
+
+ auto& dev_ctx = ctx.template device_context();
+ const auto& mkldnn_engine = dev_ctx.GetEngine();
+
+ T* input_grad_data = nullptr;
+ T* w_grad_data = nullptr;
+
+ Tensor* input_grad = ctx.Output
